Improve graphics processing and resource management (#382)
Refine graphics processing and resource handling.
This commit is contained in:
Bruno
GitHub
parent
f208b8da46
commit
732d0c4cba
|
|
@ -129,8 +129,14 @@ struct VideoManager::Impl
|
|||
if (videoSettings.use_disk_shader_cache && !stop_token.stop_requested())
|
||||
{
|
||||
const u64 program_id = m_modules.OperatingSystem().GetProgramId();
|
||||
m_gpuCore->Renderer().ReadRasterizer()->LoadDiskResources(program_id, stop_token,
|
||||
notify_disk_cache_progress);
|
||||
try {
|
||||
m_gpuCore->Renderer().ReadRasterizer()->LoadDiskResources(
|
||||
program_id, stop_token, notify_disk_cache_progress);
|
||||
} catch (const std::exception& exception) {
|
||||
// std::async stores exceptions in its future. If one escapes here,
|
||||
// Start is never called and guest GPU requests wait indefinitely.
|
||||
LOG_ERROR(Render, "Shader cache loading failed: {}", exception.what());
|
||||
}
|
||||
}
|
||||
}
|
||||
notify_disk_cache_progress(VideoCore::LoadCallbackStage::Complete, 0, 0);
|
||||
|
|
|
|||
|
|
@ -226,6 +226,7 @@ add_library(shader_recompiler STATIC
|
|||
ir_opt/lower_fp16_to_fp32.cpp
|
||||
ir_opt/lower_fp64_to_fp32.cpp
|
||||
ir_opt/lower_int64_to_int32.cpp
|
||||
ir_opt/bindless_descriptor.h
|
||||
ir_opt/passes.h
|
||||
ir_opt/position_pass.cpp
|
||||
ir_opt/rescaling_pass.cpp
|
||||
|
|
|
|||
|
|
@ -398,6 +398,13 @@ void SetupTransformFeedbackCapabilities(EmitContext& ctx, Id main_func) {
|
|||
}
|
||||
|
||||
void SetupCapabilities(const Profile& profile, const Info& info, EmitContext& ctx) {
|
||||
if (ctx.uses_nonuniform_sampled_image) {
|
||||
if (profile.supported_spirv < 0x00010500) {
|
||||
ctx.AddExtension("SPV_EXT_descriptor_indexing");
|
||||
}
|
||||
ctx.AddCapability(spv::Capability::ShaderNonUniform);
|
||||
ctx.AddCapability(spv::Capability::SampledImageArrayNonUniformIndexing);
|
||||
}
|
||||
if (info.uses_sampled_1d) {
|
||||
ctx.AddCapability(spv::Capability::Sampled1D);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -227,11 +227,26 @@ private:
|
|||
spv::ImageOperandsMask mask{};
|
||||
};
|
||||
|
||||
Id Texture(EmitContext& ctx, IR::TextureInstInfo info, [[maybe_unused]] const IR::Value& index) {
|
||||
Id Texture(EmitContext& ctx, IR::TextureInstInfo info, const IR::Value& index) {
|
||||
const TextureDefinition& def{ctx.textures.at(info.descriptor_index)};
|
||||
if (def.count > 1) {
|
||||
const Id pointer{ctx.OpAccessChain(def.pointer_type, def.id, ctx.Def(index))};
|
||||
return ctx.OpLoad(def.sampled_type, pointer);
|
||||
if (!index.IsImmediate() && !ctx.profile.support_sampled_image_array_nonuniform_indexing) {
|
||||
throw NotImplementedException("Nonuniform sampled texture indexing is unsupported by the host");
|
||||
}
|
||||
const Id array_index{ctx.Def(index)};
|
||||
const Id pointer{ctx.OpAccessChain(def.pointer_type, def.id, array_index)};
|
||||
const Id sampled_image{ctx.OpLoad(def.sampled_type, pointer)};
|
||||
// Guest bindless handles can depend on interpolated surface coordinates. Without
|
||||
// NonUniform, the host may select a single array element for the entire subgroup.
|
||||
if (!index.IsImmediate() && ctx.profile.support_sampled_image_array_nonuniform_indexing) {
|
||||
ctx.uses_nonuniform_sampled_image = true;
|
||||
for (const Id object : {array_index, pointer, sampled_image}) {
|
||||
if (ctx.nonuniform_sampled_ids.insert(object.value).second) {
|
||||
ctx.Decorate(object, spv::Decoration::NonUniform);
|
||||
}
|
||||
}
|
||||
}
|
||||
return sampled_image;
|
||||
} else {
|
||||
return ctx.OpLoad(def.sampled_type, def.id);
|
||||
}
|
||||
|
|
@ -318,8 +333,12 @@ Id IsScaled(EmitContext& ctx, const IR::Value& index, Id member_index, u32 base_
|
|||
if (base_index != 0) {
|
||||
index_value = ctx.OpIAdd(ctx.U32[1], index_value, ctx.Const(base_index));
|
||||
}
|
||||
const Id word_index{ctx.OpShiftRightLogical(ctx.U32[1], index_value, ctx.Const(5u))};
|
||||
const Id pointer{ctx.OpAccessChain(push_constant_u32, ctx.rescaling_push_constants,
|
||||
member_index, word_index)};
|
||||
const Id word{ctx.OpLoad(ctx.U32[1], pointer)};
|
||||
const Id bit_index{ctx.OpBitwiseAnd(ctx.U32[1], index_value, ctx.Const(31u))};
|
||||
bit = ctx.OpBitFieldUExtract(ctx.U32[1], index_value, bit_index, ctx.Const(1u));
|
||||
bit = ctx.OpBitFieldUExtract(ctx.U32[1], word, bit_index, ctx.Const(1u));
|
||||
}
|
||||
return ctx.OpINotEqual(ctx.U1, bit, ctx.u32_zero_value);
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1517,6 +1517,8 @@ void EmitContext::DefineInputs(const IR::Program& program) {
|
|||
}
|
||||
if (info.uses_sample_id) {
|
||||
sample_id = DefineInput(*this, U32[1], false, spv::BuiltIn::SampleId);
|
||||
// Vulkan requires Flat on integer fragment inputs, including built-ins.
|
||||
Decorate(sample_id, spv::Decoration::Flat);
|
||||
}
|
||||
if (info.uses_is_helper_invocation) {
|
||||
is_helper_invocation = DefineInput(*this, U1, false, spv::BuiltIn::HelperInvocation);
|
||||
|
|
|
|||
|
|
@ -4,6 +4,7 @@
|
|||
#pragma once
|
||||
|
||||
#include <array>
|
||||
#include <unordered_set>
|
||||
|
||||
#include <sirit/sirit.h>
|
||||
|
||||
|
|
@ -254,6 +255,9 @@ public:
|
|||
std::vector<TextureDefinition> textures;
|
||||
std::vector<ImageDefinition> images;
|
||||
|
||||
bool uses_nonuniform_sampled_image{};
|
||||
std::unordered_set<u32> nonuniform_sampled_ids;
|
||||
|
||||
Id workgroup_id{};
|
||||
Id local_invocation_id{};
|
||||
Id invocation_id{};
|
||||
|
|
|
|||
|
|
@ -20,6 +20,9 @@ public:
|
|||
|
||||
[[nodiscard]] virtual u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) = 0;
|
||||
|
||||
// Zero means this environment cannot provide the bound size (e.g. an old cache).
|
||||
[[nodiscard]] virtual u32 ReadCbufSize([[maybe_unused]] u32 cbuf_index) { return 0; }
|
||||
|
||||
[[nodiscard]] virtual TextureType ReadTextureType(u32 raw_handle) = 0;
|
||||
|
||||
[[nodiscard]] virtual TexturePixelFormat ReadTexturePixelFormat(u32 raw_handle) = 0;
|
||||
|
|
|
|||
|
|
@ -21,6 +21,7 @@ struct HostTranslateInfo {
|
|||
///< passthrough shaders
|
||||
bool support_conditional_barrier{}; ///< True when the device supports barriers in conditional
|
||||
///< control flow
|
||||
u32 max_bindless_descriptors_per_stage{8};
|
||||
};
|
||||
|
||||
} // namespace Shader
|
||||
|
|
|
|||
|
|
@ -0,0 +1,14 @@
|
|||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
#pragma once
|
||||
#include "yuzu_common/common_types.h"
|
||||
|
||||
namespace Shader::Optimization {
|
||||
// A selected handle field must fit entirely in the bound constant buffer.
|
||||
constexpr u32 BindlessDescriptorCount(u32 buffer_size, u32 base_offset, u32 size_shift) {
|
||||
if (size_shift > 31 || base_offset > buffer_size ||
|
||||
buffer_size - base_offset < sizeof(u32)) {
|
||||
return 0;
|
||||
}
|
||||
return 1 + (buffer_size - base_offset - sizeof(u32)) / (1U << size_shift);
|
||||
}
|
||||
}
|
||||
|
|
@ -1,14 +1,17 @@
|
|||
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
|
||||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
|
||||
#include <algorithm>
|
||||
#include <optional>
|
||||
#include <queue>
|
||||
#include <unordered_map>
|
||||
#include <unordered_set>
|
||||
|
||||
#include <boost/container/flat_set.hpp>
|
||||
#include <boost/container/small_vector.hpp>
|
||||
|
||||
#include "yuzu_common/alignment.h"
|
||||
#include "yuzu_shader_recompiler/frontend/ir/basic_block.h"
|
||||
#include "yuzu_shader_recompiler/frontend/ir/breadth_first_search.h"
|
||||
#include "yuzu_shader_recompiler/frontend/ir/ir_emitter.h"
|
||||
#include "yuzu_shader_recompiler/frontend/ir/value.h"
|
||||
#include "yuzu_shader_recompiler/host_translate_info.h"
|
||||
|
|
@ -329,8 +332,50 @@ std::optional<LowAddrInfo> TrackLowAddress(IR::Inst* inst) {
|
|||
};
|
||||
}
|
||||
|
||||
/// Tries to track the storage buffer address used by a global memory instruction
|
||||
std::optional<StorageBufferAddr> Track(const IR::Value& value, const Bias* bias) {
|
||||
/// Memory reads produce a value unrelated to the address used to locate it.
|
||||
bool IsIndirectPointerLoad(IR::Opcode opcode) {
|
||||
switch (opcode) {
|
||||
case IR::Opcode::LoadGlobalS8:
|
||||
case IR::Opcode::LoadGlobalU8:
|
||||
case IR::Opcode::LoadGlobalS16:
|
||||
case IR::Opcode::LoadGlobalU16:
|
||||
case IR::Opcode::LoadGlobal32:
|
||||
case IR::Opcode::LoadGlobal64:
|
||||
case IR::Opcode::LoadGlobal128:
|
||||
case IR::Opcode::LoadStorageS8:
|
||||
case IR::Opcode::LoadStorageU8:
|
||||
case IR::Opcode::LoadStorageS16:
|
||||
case IR::Opcode::LoadStorageU16:
|
||||
case IR::Opcode::LoadStorage32:
|
||||
case IR::Opcode::LoadStorage64:
|
||||
case IR::Opcode::LoadStorage128:
|
||||
case IR::Opcode::LoadSharedS8:
|
||||
case IR::Opcode::LoadSharedU8:
|
||||
case IR::Opcode::LoadSharedS16:
|
||||
case IR::Opcode::LoadSharedU16:
|
||||
case IR::Opcode::LoadSharedU32:
|
||||
case IR::Opcode::LoadSharedU64:
|
||||
case IR::Opcode::LoadSharedU128:
|
||||
case IR::Opcode::LoadLocal:
|
||||
case IR::Opcode::GetCbufU8:
|
||||
case IR::Opcode::GetCbufS8:
|
||||
case IR::Opcode::GetCbufU16:
|
||||
case IR::Opcode::GetCbufS16:
|
||||
case IR::Opcode::GetCbufF32:
|
||||
return true;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
struct TrackingState {
|
||||
std::unordered_map<const IR::Inst*, std::optional<StorageBufferAddr>> selections;
|
||||
std::unordered_set<const IR::Inst*> active;
|
||||
};
|
||||
|
||||
/// Tries to track a base without crossing pointer loads or choosing one side of a merge.
|
||||
std::optional<StorageBufferAddr> TrackImpl(const IR::Value& value, const Bias* bias,
|
||||
bool* indirect_pointer, TrackingState& state) {
|
||||
const auto pred{[bias](const IR::Inst* inst) -> std::optional<StorageBufferAddr> {
|
||||
if (inst->GetOpcode() != IR::Opcode::GetCbufU32 &&
|
||||
inst->GetOpcode() != IR::Opcode::GetCbufU32x2) {
|
||||
|
|
@ -363,7 +408,71 @@ std::optional<StorageBufferAddr> Track(const IR::Value& value, const Bias* bias)
|
|||
}
|
||||
return storage_buffer;
|
||||
}};
|
||||
return BreadthFirstSearch(value, pred);
|
||||
if (value.IsImmediate()) {
|
||||
return std::nullopt;
|
||||
}
|
||||
small_vector<const IR::Inst*, 16> visited;
|
||||
std::queue<const IR::Inst*> pending;
|
||||
const auto enqueue = [&](const IR::Value& arg) {
|
||||
if (arg.IsImmediate()) return;
|
||||
const auto* inst = arg.InstRecursive();
|
||||
if (std::ranges::find(visited, inst) == visited.end()) {
|
||||
visited.push_back(inst);
|
||||
pending.push(inst);
|
||||
}
|
||||
};
|
||||
enqueue(value);
|
||||
while (!pending.empty()) {
|
||||
const auto* inst = pending.front();
|
||||
pending.pop();
|
||||
if (const auto result = pred(inst)) return result;
|
||||
if (inst->GetOpcode() == IR::Opcode::SelectU32 || inst->GetOpcode() == IR::Opcode::Phi) {
|
||||
auto cached = state.selections.find(inst);
|
||||
if (cached == state.selections.end()) {
|
||||
// Cyclic/very deep SSA merges remain on the global-memory path.
|
||||
if (state.active.size() >= 64 || !state.active.insert(inst).second) {
|
||||
if (indirect_pointer) *indirect_pointer = true;
|
||||
continue;
|
||||
}
|
||||
std::optional<StorageBufferAddr> common;
|
||||
bool mismatch = false;
|
||||
const size_t begin = inst->GetOpcode() == IR::Opcode::SelectU32 ? 1 : 0;
|
||||
for (size_t arg = begin; arg < inst->NumArgs(); ++arg) {
|
||||
const auto branch = TrackImpl(inst->Arg(arg), bias, indirect_pointer, state);
|
||||
if (!branch || (common && *common != *branch)) {
|
||||
mismatch = true;
|
||||
break;
|
||||
}
|
||||
common = branch;
|
||||
}
|
||||
state.active.erase(inst);
|
||||
cached = state.selections.emplace(inst, mismatch ? std::nullopt : common).first;
|
||||
}
|
||||
if (cached->second) return cached->second;
|
||||
if (indirect_pointer) *indirect_pointer = true;
|
||||
// A separate additive base may still be present in the pending queue.
|
||||
continue;
|
||||
}
|
||||
if (IsIndirectPointerLoad(inst->GetOpcode())) {
|
||||
// Its address describes where the pointer is stored, not where it points.
|
||||
// Following that address can wrongly select an unrelated storage buffer.
|
||||
if (indirect_pointer) *indirect_pointer = true;
|
||||
continue;
|
||||
}
|
||||
if (inst->GetOpcode() == IR::Opcode::GetCbufU32 ||
|
||||
inst->GetOpcode() == IR::Opcode::GetCbufU32x2) {
|
||||
// A dynamically indexed descriptor cannot be identified by tracking its index.
|
||||
continue;
|
||||
}
|
||||
for (size_t arg = inst->NumArgs(); arg--;) enqueue(inst->Arg(arg));
|
||||
}
|
||||
return std::nullopt;
|
||||
}
|
||||
|
||||
std::optional<StorageBufferAddr> Track(const IR::Value& value, const Bias* bias,
|
||||
bool* indirect_pointer = nullptr) {
|
||||
TrackingState state;
|
||||
return TrackImpl(value, bias, indirect_pointer, state);
|
||||
}
|
||||
|
||||
/// Collects the storage buffer used by a global memory instruction and the instruction itself
|
||||
|
|
@ -387,10 +496,16 @@ void CollectStorageBuffers(IR::Block& block, IR::Inst& inst, StorageInfo& info)
|
|||
std::optional<StorageBufferAddr> storage_buffer{Track(low_addr, &nvn_bias)};
|
||||
if (!storage_buffer) {
|
||||
// If it fails, track without a bias
|
||||
storage_buffer = Track(low_addr, nullptr);
|
||||
bool indirect_pointer{};
|
||||
storage_buffer = Track(low_addr, nullptr, &indirect_pointer);
|
||||
if (!storage_buffer) {
|
||||
// If that also fails, use NVN fallbacks
|
||||
LOG_WARNING(Shader, "Storage buffer failed to track, using global memory fallbacks");
|
||||
if (indirect_pointer) {
|
||||
LOG_INFO(Shader, "Storage buffer pointer loaded from memory; "
|
||||
"using runtime global memory lookup");
|
||||
} else {
|
||||
LOG_WARNING(Shader, "Storage buffer failed to track, using global memory fallbacks");
|
||||
}
|
||||
return;
|
||||
}
|
||||
LOG_WARNING(Shader, "Storage buffer tracked without bias, index {} offset {}",
|
||||
|
|
|
|||
|
|
@ -4,6 +4,8 @@
|
|||
#include <algorithm>
|
||||
#include <bit>
|
||||
#include <optional>
|
||||
#include <unordered_map>
|
||||
#include <unordered_set>
|
||||
|
||||
#include <boost/container/small_vector.hpp>
|
||||
|
||||
|
|
@ -13,10 +15,14 @@
|
|||
#include "yuzu_shader_recompiler/frontend/ir/ir_emitter.h"
|
||||
#include "yuzu_shader_recompiler/host_translate_info.h"
|
||||
#include "yuzu_shader_recompiler/ir_opt/passes.h"
|
||||
#include "yuzu_shader_recompiler/ir_opt/bindless_descriptor.h"
|
||||
#include "yuzu_shader_recompiler/shader_info.h"
|
||||
|
||||
namespace Shader::Optimization {
|
||||
namespace {
|
||||
constexpr u32 DESCRIPTOR_SIZE = 8;
|
||||
constexpr u32 DESCRIPTOR_SIZE_SHIFT = static_cast<u32>(std::countr_zero(DESCRIPTOR_SIZE));
|
||||
|
||||
struct ConstBufferAddr {
|
||||
u32 index;
|
||||
u32 offset;
|
||||
|
|
@ -26,6 +32,7 @@ struct ConstBufferAddr {
|
|||
u32 secondary_shift_left;
|
||||
IR::U32 dynamic_offset;
|
||||
u32 count;
|
||||
u32 size_shift{DESCRIPTOR_SIZE_SHIFT};
|
||||
bool has_secondary;
|
||||
};
|
||||
|
||||
|
|
@ -37,8 +44,76 @@ struct TextureInst {
|
|||
|
||||
using TextureInstVector = boost::container::small_vector<TextureInst, 24>;
|
||||
|
||||
constexpr u32 DESCRIPTOR_SIZE = 8;
|
||||
constexpr u32 DESCRIPTOR_SIZE_SHIFT = static_cast<u32>(std::countr_zero(DESCRIPTOR_SIZE));
|
||||
// Bindless handles may be fields of material records rather than consecutive handles.
|
||||
// Track alignment through SSA merges so indexing and descriptor uploads use the same stride.
|
||||
u32 OffsetAlignment(const IR::Value& value, std::unordered_set<const IR::Inst*>& visiting,
|
||||
std::unordered_map<const IR::Inst*, u32>& alignments, u32 depth = 0) {
|
||||
if (value.IsImmediate()) {
|
||||
return static_cast<u32>(std::countr_zero(value.U32()));
|
||||
}
|
||||
const IR::Inst* inst{value.InstRecursive()};
|
||||
if (depth >= 128 || !inst) {
|
||||
return 0;
|
||||
}
|
||||
if (const auto cached = alignments.find(inst); cached != alignments.end()) {
|
||||
return cached->second;
|
||||
}
|
||||
if (!visiting.insert(inst).second) {
|
||||
return 0;
|
||||
}
|
||||
const auto alignment = [&] {
|
||||
switch (inst->GetOpcode()) {
|
||||
case IR::Opcode::UndefU32:
|
||||
// Undefined predecessors impose no constraint on the defined value.
|
||||
return 32U;
|
||||
case IR::Opcode::ShiftLeftLogical32:
|
||||
if (inst->Arg(1).IsImmediate() && inst->Arg(1).U32() < 32) {
|
||||
return std::min(32U, inst->Arg(1).U32() +
|
||||
OffsetAlignment(inst->Arg(0), visiting, alignments, depth + 1));
|
||||
}
|
||||
return 0U;
|
||||
case IR::Opcode::Phi: {
|
||||
u32 result{32};
|
||||
for (size_t i = 0; i < inst->NumArgs(); ++i) {
|
||||
result = std::min(result, OffsetAlignment(inst->Arg(i), visiting, alignments, depth + 1));
|
||||
}
|
||||
return result;
|
||||
}
|
||||
case IR::Opcode::BitwiseAnd32:
|
||||
// AND guarantees every zero low bit guaranteed by either operand.
|
||||
return std::max(OffsetAlignment(inst->Arg(0), visiting, alignments, depth + 1),
|
||||
OffsetAlignment(inst->Arg(1), visiting, alignments, depth + 1));
|
||||
case IR::Opcode::IAdd32:
|
||||
case IR::Opcode::ISub32:
|
||||
case IR::Opcode::BitwiseOr32:
|
||||
case IR::Opcode::BitwiseXor32:
|
||||
// Addition/subtraction wrap in U32, preserving the common power-of-two divisor.
|
||||
return std::min(OffsetAlignment(inst->Arg(0), visiting, alignments, depth + 1),
|
||||
OffsetAlignment(inst->Arg(1), visiting, alignments, depth + 1));
|
||||
case IR::Opcode::SelectU32:
|
||||
return std::min(OffsetAlignment(inst->Arg(1), visiting, alignments, depth + 1),
|
||||
OffsetAlignment(inst->Arg(2), visiting, alignments, depth + 1));
|
||||
default:
|
||||
return 0U;
|
||||
}
|
||||
}();
|
||||
visiting.erase(inst);
|
||||
// SSA is a graph: repeated phi/select predecessors must not multiply traversal work.
|
||||
alignments.emplace(inst, alignment);
|
||||
return alignment;
|
||||
}
|
||||
|
||||
u32 DynamicDescriptorSizeShift(const ConstBufferAddr& addr,
|
||||
std::unordered_set<const IR::Inst*>& visiting,
|
||||
std::unordered_map<const IR::Inst*, u32>& alignments) {
|
||||
if (addr.count <= 1) {
|
||||
return DESCRIPTOR_SIZE_SHIFT;
|
||||
}
|
||||
const u32 alignment{OffsetAlignment(addr.dynamic_offset, visiting, alignments)};
|
||||
// Alignment is a guaranteed divisor of every possible byte offset, not a minimum
|
||||
// eight-byte record size. A U32 handle can occupy a four-byte field (or be unaligned).
|
||||
return std::min(alignment, 31U);
|
||||
}
|
||||
|
||||
IR::Opcode IndexedInstruction(const IR::Inst& inst) {
|
||||
switch (inst.GetOpcode()) {
|
||||
|
|
@ -596,6 +671,33 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
to_replace.push_back(MakeInst(env, block, inst));
|
||||
}
|
||||
}
|
||||
// All source expressions are still unchanged here. Share SSA results across accesses.
|
||||
std::unordered_set<const IR::Inst*> alignment_visiting;
|
||||
std::unordered_map<const IR::Inst*, u32> offset_alignments;
|
||||
for (auto& texture : to_replace) {
|
||||
if (texture.cbuf.count <= 1) continue;
|
||||
const u32 buffer_size = env.ReadCbufSize(texture.cbuf.index);
|
||||
if (buffer_size == 0) {
|
||||
throw NotImplementedException("Bindless constant buffer size is unavailable");
|
||||
}
|
||||
const u32 shift =
|
||||
DynamicDescriptorSizeShift(texture.cbuf, alignment_visiting, offset_alignments);
|
||||
texture.cbuf.size_shift = shift;
|
||||
// Enumerate the entire buffer at this field's alignment phase. Guest U32 additions
|
||||
// can wrap and select valid handles before the original immediate base.
|
||||
const u32 phase = texture.cbuf.offset & ((1U << shift) - 1);
|
||||
const u32 adjustment = texture.cbuf.offset - phase;
|
||||
if (adjustment != 0) {
|
||||
IR::IREmitter ir{*texture.block, IR::Block::InstructionList::s_iterator_to(*texture.inst)};
|
||||
texture.cbuf.dynamic_offset = IR::U32{ir.IAdd(texture.cbuf.dynamic_offset,
|
||||
ir.Imm32(adjustment))};
|
||||
}
|
||||
texture.cbuf.offset = phase;
|
||||
texture.cbuf.count = BindlessDescriptorCount(buffer_size, phase, shift);
|
||||
if (texture.cbuf.count == 0) {
|
||||
throw NotImplementedException("Bindless constant buffer contains no complete handle");
|
||||
}
|
||||
}
|
||||
// Sort instructions to visit textures by constant buffer index, then by offset
|
||||
std::ranges::sort(to_replace, [](const auto& lhs, const auto& rhs) {
|
||||
return lhs.cbuf.offset < rhs.cbuf.offset;
|
||||
|
|
@ -615,6 +717,7 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
inst->ReplaceOpcode(IndexedInstruction(*inst));
|
||||
|
||||
const auto& cbuf{texture_inst.cbuf};
|
||||
const u32 size_shift{cbuf.size_shift};
|
||||
auto flags{inst->Flags<IR::TextureInstInfo>()};
|
||||
bool is_multisample{false};
|
||||
switch (inst->GetOpcode()) {
|
||||
|
|
@ -681,7 +784,7 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
.cbuf_index = cbuf.index,
|
||||
.cbuf_offset = cbuf.offset,
|
||||
.count = cbuf.count,
|
||||
.size_shift = DESCRIPTOR_SIZE_SHIFT,
|
||||
.size_shift = size_shift,
|
||||
});
|
||||
} else {
|
||||
index = descriptors.Add(ImageDescriptor{
|
||||
|
|
@ -693,7 +796,7 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
.cbuf_index = cbuf.index,
|
||||
.cbuf_offset = cbuf.offset,
|
||||
.count = cbuf.count,
|
||||
.size_shift = DESCRIPTOR_SIZE_SHIFT,
|
||||
.size_shift = size_shift,
|
||||
});
|
||||
}
|
||||
break;
|
||||
|
|
@ -709,7 +812,7 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
.secondary_cbuf_offset = cbuf.secondary_offset,
|
||||
.secondary_shift_left = cbuf.secondary_shift_left,
|
||||
.count = cbuf.count,
|
||||
.size_shift = DESCRIPTOR_SIZE_SHIFT,
|
||||
.size_shift = size_shift,
|
||||
});
|
||||
} else {
|
||||
index = descriptors.Add(TextureDescriptor{
|
||||
|
|
@ -724,7 +827,7 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
.secondary_cbuf_offset = cbuf.secondary_offset,
|
||||
.secondary_shift_left = cbuf.secondary_shift_left,
|
||||
.count = cbuf.count,
|
||||
.size_shift = DESCRIPTOR_SIZE_SHIFT,
|
||||
.size_shift = size_shift,
|
||||
});
|
||||
}
|
||||
break;
|
||||
|
|
@ -735,9 +838,9 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
if (cbuf.count > 1) {
|
||||
const auto insert_point{IR::Block::InstructionList::s_iterator_to(*inst)};
|
||||
IR::IREmitter ir{*texture_inst.block, insert_point};
|
||||
const IR::U32 shift{ir.Imm32(std::countr_zero(DESCRIPTOR_SIZE))};
|
||||
inst->SetArg(0, ir.UMin(ir.ShiftRightArithmetic(cbuf.dynamic_offset, shift),
|
||||
ir.Imm32(DESCRIPTOR_SIZE - 1)));
|
||||
const IR::U32 shift{ir.Imm32(size_shift)};
|
||||
inst->SetArg(0, ir.UMin(ir.ShiftRightLogical(cbuf.dynamic_offset, shift),
|
||||
ir.Imm32(cbuf.count - 1)));
|
||||
} else {
|
||||
inst->SetArg(0, IR::Value{});
|
||||
}
|
||||
|
|
@ -750,6 +853,15 @@ void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo
|
|||
}
|
||||
}
|
||||
}
|
||||
// Count actual bindings after deduplication, rather than texture instructions.
|
||||
// Never map a valid guest index to another handle to satisfy host resource limits.
|
||||
const u64 sampled_count = u64{NumDescriptors(program.info.texture_descriptors)} +
|
||||
NumDescriptors(program.info.texture_buffer_descriptors);
|
||||
if (std::ranges::any_of(to_replace, [](const auto& texture) {
|
||||
return !texture.cbuf.dynamic_offset.IsEmpty();
|
||||
}) && sampled_count > host_info.max_bindless_descriptors_per_stage) {
|
||||
throw NotImplementedException("Sampled texture bindings exceed the host descriptor budget");
|
||||
}
|
||||
}
|
||||
|
||||
void JoinTextureInfo(Info& base, Info& source) {
|
||||
|
|
|
|||
|
|
@ -11,6 +11,7 @@ struct Profile {
|
|||
u32 supported_spirv{0x00010000};
|
||||
bool unified_descriptor_binding{};
|
||||
bool support_descriptor_aliasing{};
|
||||
bool support_sampled_image_array_nonuniform_indexing{};
|
||||
// MoltenVK needs one canonical uniform view; storage-buffer aliases remain available.
|
||||
bool disable_uniform_buffer_descriptor_aliasing{};
|
||||
bool support_int8{};
|
||||
|
|
|
|||
|
|
@ -42,6 +42,7 @@
|
|||
<ClInclude Include="backend\spirv\emit_spirv.h" />
|
||||
<ClInclude Include="backend\spirv\emit_spirv_instructions.h" />
|
||||
<ClInclude Include="backend\spirv\spirv_emit_context.h" />
|
||||
<ClInclude Include="ir_opt\bindless_descriptor.h" />
|
||||
<ClInclude Include="environment.h" />
|
||||
<ClInclude Include="exception.h" />
|
||||
<ClInclude Include="frontend\ir\abstract_syntax_list.h" />
|
||||
|
|
|
|||
|
|
@ -1,4 +1,4 @@
|
|||
<?xml version="1.0" encoding="utf-8"?>
|
||||
<?xml version="1.0" encoding="utf-8"?>
|
||||
<Project ToolsVersion="4.0" xmlns="http://schemas.microsoft.com/developer/msbuild/2003">
|
||||
<ItemGroup>
|
||||
<Filter Include="Source Files">
|
||||
|
|
@ -237,7 +237,12 @@
|
|||
<Filter>Header Files\frontend\maxwell\translate\impl</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="ir_opt\passes.h">
|
||||
<Filter>Header Files\ir_opt</Filter>
|
||||
<Filter>Header Files</Filter>
|
||||
</ClInclude>
|
||||
</ItemGroup>
|
||||
<ItemGroup>
|
||||
<ClInclude Include="ir_opt\bindless_descriptor.h">
|
||||
<Filter>Header Files</Filter>
|
||||
</ClInclude>
|
||||
</ItemGroup>
|
||||
<ItemGroup>
|
||||
|
|
|
|||
|
|
@ -197,6 +197,10 @@ add_library(video_core STATIC
|
|||
renderer_vulkan/present/window_adapt_pass.h
|
||||
renderer_vulkan/blit_image.cpp
|
||||
renderer_vulkan/blit_image.h
|
||||
renderer_vulkan/descriptor_lookup.h
|
||||
renderer_vulkan/descriptor_limits.h
|
||||
renderer_vulkan/descriptor_payload.h
|
||||
renderer_vulkan/cbuf_size_dependencies.h
|
||||
renderer_vulkan/fixed_pipeline_state.cpp
|
||||
renderer_vulkan/fixed_pipeline_state.h
|
||||
renderer_vulkan/maxwell_to_vk.cpp
|
||||
|
|
@ -277,6 +281,7 @@ add_library(video_core STATIC
|
|||
texture_cache/image_base.h
|
||||
texture_cache/image_info.cpp
|
||||
texture_cache/image_info.h
|
||||
texture_cache/mip_levels.h
|
||||
texture_cache/image_view_base.cpp
|
||||
texture_cache/image_view_base.h
|
||||
texture_cache/image_view_info.cpp
|
||||
|
|
|
|||
|
|
@ -0,0 +1,93 @@
|
|||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <array>
|
||||
#include <bit>
|
||||
#include <limits>
|
||||
#include <unordered_map>
|
||||
|
||||
#include "yuzu_common/common_types.h"
|
||||
|
||||
namespace Vulkan {
|
||||
|
||||
// Updated by the GPU thread and disk-cache reader, never by pipeline worker jobs.
|
||||
class CbufSizeDependencies {
|
||||
public:
|
||||
void Observe(u64 shader, u32 mask) {
|
||||
if (shader != 0) {
|
||||
const auto [it, inserted] = masks.try_emplace(shader, mask);
|
||||
const u32 combined = it->second | mask;
|
||||
if (inserted || combined != it->second) {
|
||||
it->second = combined;
|
||||
++revision;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
u32 Mask(u64 shader) const {
|
||||
if (shader == 0) {
|
||||
return 0;
|
||||
}
|
||||
const auto it = masks.find(shader);
|
||||
// Until translation has reported its dependencies, specialize conservatively.
|
||||
return it != masks.end() ? it->second : std::numeric_limits<u32>::max();
|
||||
}
|
||||
|
||||
u64 Revision() const noexcept { return revision; }
|
||||
|
||||
template <size_t NumBanks>
|
||||
static void Normalize(std::array<u32, NumBanks>& sizes, u32 mask) {
|
||||
static_assert(NumBanks <= 32);
|
||||
for (size_t bank = 0; bank < NumBanks; ++bank) {
|
||||
if ((mask & (1U << bank)) == 0) {
|
||||
sizes[bank] = 0;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
private:
|
||||
std::unordered_map<u64, u32> masks;
|
||||
u64 revision{};
|
||||
};
|
||||
|
||||
// Retain only dependency metadata, never bound sizes or guest addresses. Required banks
|
||||
// are read on every draw, including enable/disable and size changes with the same shader.
|
||||
template <size_t NumStages, size_t NumBanks>
|
||||
class GraphicsCbufSizeState {
|
||||
public:
|
||||
GraphicsCbufSizeState() { stage_masks.fill(std::numeric_limits<u32>::max()); }
|
||||
|
||||
template <typename ReadSize>
|
||||
void Refresh(std::array<std::array<u32, NumBanks>, NumStages>& sizes,
|
||||
const std::array<u64, NumStages + 1>& hashes,
|
||||
const CbufSizeDependencies& dependencies, ReadSize&& read_size) {
|
||||
static_assert(NumBanks < 32);
|
||||
if (revision != dependencies.Revision() || shader_hashes != hashes) {
|
||||
shader_hashes = hashes;
|
||||
revision = dependencies.Revision();
|
||||
for (size_t stage = 0; stage < NumStages; ++stage) {
|
||||
const u32 mask = (dependencies.Mask(hashes[stage + 1]) |
|
||||
(stage == 0 ? dependencies.Mask(hashes[0]) : 0)) & ((1U << NumBanks) - 1);
|
||||
if (mask != stage_masks[stage]) {
|
||||
sizes[stage].fill(0);
|
||||
stage_masks[stage] = mask;
|
||||
}
|
||||
}
|
||||
}
|
||||
for (size_t stage = 0; stage < NumStages; ++stage) {
|
||||
u32 mask = stage_masks[stage];
|
||||
while (mask != 0) {
|
||||
const auto bank = std::countr_zero(mask);
|
||||
sizes[stage][bank] = read_size(stage, bank);
|
||||
mask &= mask - 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
private:
|
||||
std::array<u64, NumStages + 1> shader_hashes{};
|
||||
std::array<u32, NumStages> stage_masks{};
|
||||
u64 revision{std::numeric_limits<u64>::max()};
|
||||
};
|
||||
|
||||
} // namespace Vulkan
|
||||
|
|
@ -0,0 +1,99 @@
|
|||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <array>
|
||||
#include <cstddef>
|
||||
#include <cstdint>
|
||||
#include <span>
|
||||
|
||||
#include <vulkan/vulkan_core.h>
|
||||
|
||||
namespace Vulkan {
|
||||
namespace DescriptorLimitsDetail {
|
||||
struct DescriptorUsage {
|
||||
uint64_t samplers{}, sampled_images{}, storage_images{}, uniform_buffers{}, storage_buffers{};
|
||||
uint64_t resources{};
|
||||
|
||||
bool Add(VkDescriptorType type, uint32_t count) {
|
||||
switch (type) {
|
||||
case VK_DESCRIPTOR_TYPE_UNIFORM_BUFFER:
|
||||
uniform_buffers += count;
|
||||
break;
|
||||
case VK_DESCRIPTOR_TYPE_STORAGE_BUFFER:
|
||||
storage_buffers += count;
|
||||
break;
|
||||
case VK_DESCRIPTOR_TYPE_UNIFORM_TEXEL_BUFFER:
|
||||
sampled_images += count;
|
||||
break;
|
||||
case VK_DESCRIPTOR_TYPE_STORAGE_TEXEL_BUFFER:
|
||||
case VK_DESCRIPTOR_TYPE_STORAGE_IMAGE:
|
||||
storage_images += count;
|
||||
break;
|
||||
case VK_DESCRIPTOR_TYPE_COMBINED_IMAGE_SAMPLER:
|
||||
samplers += count;
|
||||
sampled_images += count;
|
||||
break;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
// Combined image samplers count as one resource, but against both type limits.
|
||||
resources += count;
|
||||
return true;
|
||||
}
|
||||
};
|
||||
} // namespace DescriptorLimitsDetail
|
||||
|
||||
// Check the actual pipeline layout once at construction, including widened stage visibility.
|
||||
// No per-draw cost, fixed stage count, and 64-bit sums for untrusted descriptor counts.
|
||||
inline const char* DescriptorLimitViolation(
|
||||
const VkPhysicalDeviceLimits& limits, std::span<const VkDescriptorSetLayoutBinding> bindings,
|
||||
uint32_t color_attachments = 0) {
|
||||
constexpr std::array stages{
|
||||
VK_SHADER_STAGE_VERTEX_BIT,
|
||||
VK_SHADER_STAGE_TESSELLATION_CONTROL_BIT,
|
||||
VK_SHADER_STAGE_TESSELLATION_EVALUATION_BIT,
|
||||
VK_SHADER_STAGE_GEOMETRY_BIT,
|
||||
VK_SHADER_STAGE_FRAGMENT_BIT,
|
||||
VK_SHADER_STAGE_COMPUTE_BIT,
|
||||
};
|
||||
DescriptorLimitsDetail::DescriptorUsage total;
|
||||
std::array<DescriptorLimitsDetail::DescriptorUsage, stages.size()> usage{};
|
||||
for (const auto& binding : bindings) {
|
||||
if (!total.Add(binding.descriptorType, binding.descriptorCount)) {
|
||||
return "Unsupported descriptor type in pipeline layout";
|
||||
}
|
||||
for (std::size_t stage = 0; stage < stages.size(); ++stage) {
|
||||
if ((binding.stageFlags & stages[stage]) != 0) {
|
||||
usage[stage].Add(binding.descriptorType, binding.descriptorCount);
|
||||
}
|
||||
}
|
||||
}
|
||||
if (total.samplers > limits.maxDescriptorSetSamplers ||
|
||||
total.sampled_images > limits.maxDescriptorSetSampledImages ||
|
||||
total.storage_images > limits.maxDescriptorSetStorageImages ||
|
||||
total.uniform_buffers > limits.maxDescriptorSetUniformBuffers ||
|
||||
total.storage_buffers > limits.maxDescriptorSetStorageBuffers) {
|
||||
return "Pipeline exceeds descriptor set resource limits";
|
||||
}
|
||||
for (std::size_t stage = 0; stage < stages.size(); ++stage) {
|
||||
const auto& count = usage[stage];
|
||||
const uint64_t attachments = stages[stage] == VK_SHADER_STAGE_FRAGMENT_BIT
|
||||
? color_attachments : 0;
|
||||
if (count.samplers > limits.maxPerStageDescriptorSamplers ||
|
||||
count.sampled_images > limits.maxPerStageDescriptorSampledImages ||
|
||||
count.storage_images > limits.maxPerStageDescriptorStorageImages ||
|
||||
count.uniform_buffers > limits.maxPerStageDescriptorUniformBuffers ||
|
||||
count.storage_buffers > limits.maxPerStageDescriptorStorageBuffers ||
|
||||
count.resources + attachments > limits.maxPerStageResources) {
|
||||
return "Pipeline exceeds per-stage resource limits";
|
||||
}
|
||||
if (stages[stage] == VK_SHADER_STAGE_FRAGMENT_BIT &&
|
||||
attachments + count.storage_images + count.storage_buffers >
|
||||
limits.maxFragmentCombinedOutputResources) {
|
||||
return "Pipeline exceeds combined fragment output resource limit";
|
||||
}
|
||||
}
|
||||
return nullptr;
|
||||
}
|
||||
} // namespace Vulkan
|
||||
|
|
@ -0,0 +1,109 @@
|
|||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <array>
|
||||
#include <cstring>
|
||||
#include <limits>
|
||||
#include <utility>
|
||||
|
||||
#include "yuzu_common/common_types.h"
|
||||
|
||||
namespace Vulkan {
|
||||
|
||||
// Lives for one shader stage of one draw/dispatch. Never retains guest pointers across draws,
|
||||
// rebindings or address-space changes. Discontiguous ranges keep the original scalar read path.
|
||||
template <size_t NumBanks>
|
||||
class DescriptorCbufReader {
|
||||
public:
|
||||
template <typename Memory>
|
||||
struct BoundBuffer {
|
||||
Memory* memory;
|
||||
GPUVAddr address;
|
||||
u32 size;
|
||||
const u8* span;
|
||||
u32 Read(u32 offset) const {
|
||||
if (span && offset <= size - sizeof(u32)) {
|
||||
u32 value;
|
||||
std::memcpy(&value, span + offset, sizeof(value));
|
||||
return value;
|
||||
}
|
||||
return memory->template Read<u32>(address + offset);
|
||||
}
|
||||
};
|
||||
|
||||
// Prepare once per descriptor array. A fragmented bank takes the scalar path directly,
|
||||
// without repeating bank/size/span queries for every array element.
|
||||
template <typename Memory>
|
||||
BoundBuffer<Memory> Bind(Memory& memory, u32 bank, GPUVAddr address, u32 size, u32 count) {
|
||||
const u8* span{};
|
||||
if (count > 64 && bank < NumBanks && size >= sizeof(u32)) {
|
||||
const u32 bit = 1U << bank;
|
||||
if ((queried_banks & bit) == 0) {
|
||||
spans[bank] = memory.GetSpan(address, size);
|
||||
queried_banks |= bit;
|
||||
}
|
||||
span = spans[bank];
|
||||
}
|
||||
return {&memory, address, size, span};
|
||||
}
|
||||
|
||||
template <typename Memory>
|
||||
u32 Read(Memory& memory, u32 bank, GPUVAddr address, u32 size, u32 offset, u32 count) {
|
||||
return Bind(memory, bank, address, size, count).Read(offset);
|
||||
}
|
||||
|
||||
private:
|
||||
static_assert(NumBanks <= 32);
|
||||
std::array<const u8*, NumBanks> spans; // Only read after its queried bit is set.
|
||||
u32 queried_banks{};
|
||||
};
|
||||
|
||||
// Reset after synchronizing descriptor tables, before each draw/dispatch. Caches lookups only
|
||||
// within that operation, so pool modifications and sampler rebindings always remain visible.
|
||||
template <typename Value, size_t NumSlots = 256>
|
||||
class DrawDescriptorCache {
|
||||
public:
|
||||
DrawDescriptorCache() { Reset(); }
|
||||
|
||||
void Reset() {
|
||||
consecutive_misses = 0;
|
||||
enabled = cooldown == 0;
|
||||
if (!enabled) {
|
||||
--cooldown;
|
||||
return;
|
||||
}
|
||||
indices.fill(std::numeric_limits<u32>::max());
|
||||
}
|
||||
|
||||
bool Enabled() const noexcept { return enabled; }
|
||||
|
||||
template <typename Lookup>
|
||||
Value Get(u32 index, Lookup&& lookup) {
|
||||
if (!enabled || index == std::numeric_limits<u32>::max()) {
|
||||
return lookup(index);
|
||||
}
|
||||
const size_t slot = index & (NumSlots - 1);
|
||||
if (indices[slot] != index) {
|
||||
values[slot] = lookup(index);
|
||||
indices[slot] = index;
|
||||
if (++consecutive_misses >= 64) {
|
||||
// Correctness is unchanged: bypass useless caching and probe again later.
|
||||
enabled = false;
|
||||
cooldown = 16;
|
||||
}
|
||||
} else {
|
||||
consecutive_misses = 0;
|
||||
}
|
||||
return values[slot];
|
||||
}
|
||||
|
||||
private:
|
||||
static_assert(NumSlots != 0 && (NumSlots & (NumSlots - 1)) == 0);
|
||||
u32 consecutive_misses{}, cooldown{};
|
||||
bool enabled{true};
|
||||
std::array<u32, NumSlots> indices{};
|
||||
std::array<Value, NumSlots> values{};
|
||||
};
|
||||
|
||||
} // namespace Vulkan
|
||||
|
|
@ -0,0 +1,60 @@
|
|||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
#pragma once
|
||||
#include <array>
|
||||
#include <cstddef>
|
||||
#include <stdexcept>
|
||||
#include <vector>
|
||||
|
||||
namespace Vulkan {
|
||||
// CPU descriptor data is read by the scheduler worker after the draw is recorded.
|
||||
// Keep pending uploads stable, and drain that worker before recycling their storage.
|
||||
template <typename Entry, size_t FrameCapacity, size_t Frames = 8>
|
||||
class DescriptorPayload {
|
||||
public:
|
||||
void TickFrame() {
|
||||
frame = (frame + 1) % Frames;
|
||||
used = 0;
|
||||
overflow_active = false;
|
||||
}
|
||||
|
||||
template <typename WaitWorker>
|
||||
void Acquire(size_t count, WaitWorker&& wait_worker) {
|
||||
if (overflow_active) {
|
||||
// The ring was drained when the previous overflow upload began.
|
||||
used = 0;
|
||||
overflow_active = false;
|
||||
}
|
||||
if (count > FrameCapacity) {
|
||||
wait_worker();
|
||||
overflow.resize(count);
|
||||
upload = overflow.data();
|
||||
overflow_active = true;
|
||||
} else {
|
||||
if (count > FrameCapacity - used) {
|
||||
wait_worker();
|
||||
used = 0;
|
||||
}
|
||||
upload = ring.data() + frame * FrameCapacity + used;
|
||||
used += count;
|
||||
}
|
||||
written = 0;
|
||||
reserved = count;
|
||||
}
|
||||
|
||||
void Append(const Entry& value) {
|
||||
if (written >= reserved) {
|
||||
throw std::length_error("Descriptor upload exceeded its reservation");
|
||||
}
|
||||
upload[written++] = value;
|
||||
}
|
||||
const Entry* Data() const noexcept { return upload; }
|
||||
|
||||
private:
|
||||
static_assert(FrameCapacity > 0 && Frames > 0);
|
||||
std::array<Entry, FrameCapacity * Frames> ring;
|
||||
std::vector<Entry> overflow;
|
||||
Entry* upload{};
|
||||
size_t frame{}, used{}, reserved{}, written{};
|
||||
bool overflow_active{};
|
||||
};
|
||||
} // namespace Vulkan
|
||||
|
|
@ -10,7 +10,9 @@
|
|||
|
||||
#include "yuzu_common/common_types.h"
|
||||
#include "yuzu_shader_recompiler/backend/spirv/emit_spirv.h"
|
||||
#include "yuzu_shader_recompiler/exception.h"
|
||||
#include "yuzu_shader_recompiler/shader_info.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/descriptor_limits.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_texture_cache.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_update_descriptor.h"
|
||||
#include "yuzu_video_core/texture_cache/types.h"
|
||||
|
|
@ -55,10 +57,21 @@ inline bool RasterizesLines(Shader::OutputTopology output,
|
|||
polygon_mode == Tegra::Engines::Maxwell3D::Regs::PolygonMode::Line);
|
||||
}
|
||||
|
||||
inline size_t NumDescriptorUpdateEntries(const Shader::Info& info) {
|
||||
return size_t{Shader::NumDescriptors(info.constant_buffer_descriptors)} +
|
||||
Shader::NumDescriptors(info.storage_buffers_descriptors) +
|
||||
Shader::NumDescriptors(info.texture_buffer_descriptors) +
|
||||
Shader::NumDescriptors(info.image_buffer_descriptors) +
|
||||
Shader::NumDescriptors(info.texture_descriptors) +
|
||||
Shader::NumDescriptors(info.image_descriptors);
|
||||
}
|
||||
|
||||
class DescriptorLayoutBuilder {
|
||||
public:
|
||||
DescriptorLayoutBuilder(const Device& device_, VkShaderStageFlags storage_buffer_stages_ = 0)
|
||||
: device{&device_}, storage_buffer_stages{storage_buffer_stages_} {}
|
||||
DescriptorLayoutBuilder(const Device& device_, VkShaderStageFlags storage_buffer_stages_ = 0,
|
||||
u32 color_attachments_ = 0)
|
||||
: device{&device_}, storage_buffer_stages{storage_buffer_stages_},
|
||||
color_attachments{color_attachments_} {}
|
||||
|
||||
bool CanUsePushDescriptor() const noexcept {
|
||||
return device->IsKhrPushDescriptorSupported() &&
|
||||
|
|
@ -66,6 +79,11 @@ public:
|
|||
}
|
||||
|
||||
vk::DescriptorSetLayout CreateDescriptorSetLayout(bool use_push_descriptor) const {
|
||||
if (const auto* violation = DescriptorLimitViolation(
|
||||
device->DescriptorLimits(), std::span{bindings.data(), bindings.size()},
|
||||
color_attachments)) {
|
||||
throw Shader::NotImplementedException("{}", violation);
|
||||
}
|
||||
if (bindings.empty()) {
|
||||
return nullptr;
|
||||
}
|
||||
|
|
@ -161,12 +179,13 @@ private:
|
|||
});
|
||||
++binding;
|
||||
num_descriptors += descriptors[i].count;
|
||||
offset += sizeof(DescriptorUpdateEntry);
|
||||
offset += sizeof(DescriptorUpdateEntry) * descriptors[i].count;
|
||||
}
|
||||
}
|
||||
|
||||
const Device* device{};
|
||||
VkShaderStageFlags storage_buffer_stages{};
|
||||
u32 color_attachments{};
|
||||
bool is_compute{};
|
||||
boost::container::small_vector<VkDescriptorSetLayoutBinding, 32> bindings;
|
||||
boost::container::small_vector<VkDescriptorUpdateTemplateEntry, 32> entries;
|
||||
|
|
@ -225,6 +244,7 @@ inline void PushImageDescriptors(TextureCache& texture_cache,
|
|||
views += num_texture_buffers;
|
||||
views += num_image_buffers;
|
||||
for (const auto& desc : info.texture_descriptors) {
|
||||
bool is_rescaled{};
|
||||
for (u32 index = 0; index < desc.count; ++index) {
|
||||
const VideoCommon::ImageViewId image_view_id{(views++)->id};
|
||||
const VideoCommon::SamplerId sampler_id{*(samplers++)};
|
||||
|
|
@ -236,10 +256,13 @@ inline void PushImageDescriptors(TextureCache& texture_cache,
|
|||
const VkSampler vk_sampler{use_fallback_sampler ? sampler.HandleWithDefaultAnisotropy()
|
||||
: sampler.Handle()};
|
||||
guest_descriptor_queue.AddSampledImage(vk_image_view, vk_sampler);
|
||||
rescaling.PushTexture(texture_cache.IsRescaling(image_view));
|
||||
is_rescaled |= texture_cache.IsRescaling(image_view);
|
||||
}
|
||||
// SPIR-V allocates one scaling bit per descriptor, including array descriptors.
|
||||
rescaling.PushTexture(is_rescaled);
|
||||
}
|
||||
for (const auto& desc : info.image_descriptors) {
|
||||
bool is_rescaled{};
|
||||
for (u32 index = 0; index < desc.count; ++index) {
|
||||
ImageView& image_view{texture_cache.GetImageView((views++)->id)};
|
||||
if (desc.is_written) {
|
||||
|
|
@ -247,8 +270,9 @@ inline void PushImageDescriptors(TextureCache& texture_cache,
|
|||
}
|
||||
const VkImageView vk_image_view{image_view.StorageView(desc.type, desc.format)};
|
||||
guest_descriptor_queue.AddImage(vk_image_view);
|
||||
rescaling.PushImage(texture_cache.IsRescaling(image_view));
|
||||
is_rescaled |= texture_cache.IsRescaling(image_view);
|
||||
}
|
||||
rescaling.PushImage(is_rescaled);
|
||||
}
|
||||
}
|
||||
|
||||
|
|
|
|||
|
|
@ -318,7 +318,7 @@ std::pair<VkBuffer, VkDeviceSize> Uint8Pass::Assemble(u32 num_vertices, VkBuffer
|
|||
const u32 staging_size = static_cast<u32>(num_vertices * sizeof(u16));
|
||||
const auto staging = staging_buffer_pool.Request(staging_size, MemoryUsage::DeviceLocal);
|
||||
|
||||
compute_pass_descriptor_queue.Acquire();
|
||||
compute_pass_descriptor_queue.Acquire(2);
|
||||
compute_pass_descriptor_queue.AddBuffer(src_buffer, src_offset, num_vertices);
|
||||
compute_pass_descriptor_queue.AddBuffer(staging.buffer, staging.offset, staging_size);
|
||||
const void* const descriptor_data{compute_pass_descriptor_queue.UpdateData()};
|
||||
|
|
@ -376,7 +376,7 @@ std::pair<VkBuffer, VkDeviceSize> QuadIndexedPass::Assemble(
|
|||
const std::size_t staging_size = num_tri_vertices * sizeof(u32);
|
||||
const auto staging = staging_buffer_pool.Request(staging_size, MemoryUsage::DeviceLocal);
|
||||
|
||||
compute_pass_descriptor_queue.Acquire();
|
||||
compute_pass_descriptor_queue.Acquire(2);
|
||||
compute_pass_descriptor_queue.AddBuffer(src_buffer, src_offset, input_size);
|
||||
compute_pass_descriptor_queue.AddBuffer(staging.buffer, staging.offset, staging_size);
|
||||
const void* const descriptor_data{compute_pass_descriptor_queue.UpdateData()};
|
||||
|
|
@ -417,7 +417,7 @@ void ConditionalRenderingResolvePass::Resolve(VkBuffer dst_buffer, VkBuffer src_
|
|||
u32 src_offset, bool compare_to_zero) {
|
||||
const size_t compare_size = compare_to_zero ? 8 : 24;
|
||||
|
||||
compute_pass_descriptor_queue.Acquire();
|
||||
compute_pass_descriptor_queue.Acquire(2);
|
||||
compute_pass_descriptor_queue.AddBuffer(src_buffer, src_offset, compare_size);
|
||||
compute_pass_descriptor_queue.AddBuffer(dst_buffer, 0, sizeof(u32));
|
||||
const void* const descriptor_data{compute_pass_descriptor_queue.UpdateData()};
|
||||
|
|
@ -473,7 +473,7 @@ void QueriesPrefixScanPass::Run(VkBuffer accumulation_buffer, VkBuffer dst_buffe
|
|||
static constexpr size_t DISPATCH_SIZE = 2048U;
|
||||
size_t runs_to_do = std::min<size_t>(current_runs, DISPATCH_SIZE);
|
||||
current_runs -= runs_to_do;
|
||||
compute_pass_descriptor_queue.Acquire();
|
||||
compute_pass_descriptor_queue.Acquire(2);
|
||||
compute_pass_descriptor_queue.AddBuffer(src_buffer, 0, number_of_sums * sizeof(u64));
|
||||
compute_pass_descriptor_queue.AddBuffer(dst_buffer, 0, number_of_sums * sizeof(u64));
|
||||
compute_pass_descriptor_queue.AddBuffer(accumulation_buffer, 0, sizeof(u64));
|
||||
|
|
@ -580,7 +580,7 @@ void ASTCDecoderPass::Assemble(Image& image, const StagingBufferRef& map,
|
|||
const u32 num_dispatches_y = Common::DivCeil(swizzle.num_tiles.height, 8U);
|
||||
const u32 num_dispatches_z = image.info.resources.layers;
|
||||
|
||||
compute_pass_descriptor_queue.Acquire();
|
||||
compute_pass_descriptor_queue.Acquire(2);
|
||||
compute_pass_descriptor_queue.AddBuffer(map.buffer, input_offset,
|
||||
image.guest_size_bytes - swizzle.buffer_offset);
|
||||
compute_pass_descriptor_queue.AddImage(image.StorageImageView(swizzle.level));
|
||||
|
|
@ -685,7 +685,7 @@ void MSAACopyPass::CopyImage(Image& dst_image, Image& src_image,
|
|||
ASSERT(copy.dst_subresource.base_layer == 0);
|
||||
ASSERT(copy.dst_subresource.num_layers == 1);
|
||||
|
||||
compute_pass_descriptor_queue.Acquire();
|
||||
compute_pass_descriptor_queue.Acquire(2);
|
||||
compute_pass_descriptor_queue.AddImage(
|
||||
src_image.StorageImageView(copy.src_subresource.base_level));
|
||||
compute_pass_descriptor_queue.AddImage(
|
||||
|
|
|
|||
|
|
@ -4,8 +4,6 @@
|
|||
#include <algorithm>
|
||||
#include <vector>
|
||||
|
||||
#include <boost/container/small_vector.hpp>
|
||||
|
||||
#include "yuzu_video_core/renderer_vulkan/pipeline_helper.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/pipeline_statistics.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_buffer_cache.h"
|
||||
|
|
@ -34,6 +32,12 @@ ComputePipeline::ComputePipeline(const Device& device_, vk::PipelineCache& pipel
|
|||
: device{device_},
|
||||
pipeline_cache(pipeline_cache_), guest_descriptor_queue{guest_descriptor_queue_}, info{info_},
|
||||
spv_module(std::move(spv_module_)) {
|
||||
num_descriptor_entries = NumDescriptorUpdateEntries(info);
|
||||
num_textures = Shader::NumDescriptors(info.texture_descriptors);
|
||||
descriptor_views.resize(num_textures + Shader::NumDescriptors(info.texture_buffer_descriptors) +
|
||||
Shader::NumDescriptors(info.image_buffer_descriptors) +
|
||||
Shader::NumDescriptors(info.image_descriptors));
|
||||
descriptor_samplers.resize(num_textures);
|
||||
if (shader_notify) {
|
||||
shader_notify->MarkShaderBuilding();
|
||||
}
|
||||
|
|
@ -104,7 +108,7 @@ ComputePipeline::ComputePipeline(const Device& device_, vk::PipelineCache& pipel
|
|||
void ComputePipeline::Configure(Tegra::Engines::KeplerCompute& kepler_compute,
|
||||
Tegra::MemoryManager& gpu_memory, Scheduler& scheduler,
|
||||
BufferCache& buffer_cache, TextureCache& texture_cache) {
|
||||
guest_descriptor_queue.Acquire();
|
||||
guest_descriptor_queue.Acquire(num_descriptor_entries);
|
||||
|
||||
buffer_cache.SetComputeUniformBufferState(info.constant_buffer_mask, &uniform_buffer_sizes);
|
||||
buffer_cache.UnbindComputeStorageBuffers();
|
||||
|
|
@ -118,40 +122,61 @@ void ComputePipeline::Configure(Tegra::Engines::KeplerCompute& kepler_compute,
|
|||
|
||||
texture_cache.SynchronizeComputeDescriptors();
|
||||
|
||||
static constexpr size_t max_elements = 64;
|
||||
boost::container::static_vector<VideoCommon::ImageViewInOut, max_elements> views;
|
||||
boost::container::static_vector<VideoCommon::SamplerId, max_elements> samplers;
|
||||
auto& views = descriptor_views;
|
||||
auto& samplers = descriptor_samplers;
|
||||
size_t view_index{};
|
||||
size_t sampler_index{};
|
||||
if (num_textures > 64) {
|
||||
if (!sampler_lookup) {
|
||||
sampler_lookup = std::make_unique<DrawDescriptorCache<VideoCommon::SamplerId>>();
|
||||
}
|
||||
sampler_lookup->Reset();
|
||||
}
|
||||
const bool memoize_samplers = sampler_lookup && sampler_lookup->Enabled();
|
||||
const auto lookup_sampler = [&](u32 index) {
|
||||
return texture_cache.GetComputeSamplerId(index);
|
||||
};
|
||||
|
||||
const auto& qmd{kepler_compute.launch_description};
|
||||
const auto& cbufs{qmd.const_buffer_config};
|
||||
DescriptorCbufReader<8> cbuf_reader;
|
||||
const auto bind_cbuf = [&](u32 bank, u32 count) {
|
||||
return cbuf_reader.Bind(gpu_memory, bank, cbufs[bank].Address(), cbufs[bank].size.Value(), count);
|
||||
};
|
||||
const bool via_header_index{qmd.linked_tsc != 0};
|
||||
const auto read_handle{[&](const auto& desc, u32 index) {
|
||||
const auto prepare_handle{[&](const auto& desc) {
|
||||
ASSERT(((qmd.const_buffer_enable_mask >> desc.cbuf_index) & 1) != 0);
|
||||
const u32 index_offset{index << desc.size_shift};
|
||||
const u32 offset{desc.cbuf_offset + index_offset};
|
||||
const GPUVAddr addr{cbufs[desc.cbuf_index].Address() + offset};
|
||||
if constexpr (std::is_same_v<decltype(desc), const Shader::TextureDescriptor&> ||
|
||||
std::is_same_v<decltype(desc), const Shader::TextureBufferDescriptor&>) {
|
||||
const auto primary = bind_cbuf(desc.cbuf_index, desc.count);
|
||||
auto secondary = primary;
|
||||
bool has_secondary{};
|
||||
u32 secondary_offset{}, primary_shift{}, secondary_shift{};
|
||||
if constexpr (requires { desc.has_secondary; }) {
|
||||
if (desc.has_secondary) {
|
||||
ASSERT(((qmd.const_buffer_enable_mask >> desc.secondary_cbuf_index) & 1) != 0);
|
||||
const u32 secondary_offset{desc.secondary_cbuf_offset + index_offset};
|
||||
const GPUVAddr separate_addr{cbufs[desc.secondary_cbuf_index].Address() +
|
||||
secondary_offset};
|
||||
const u32 lhs_raw{gpu_memory.Read<u32>(addr) << desc.shift_left};
|
||||
const u32 rhs_raw{gpu_memory.Read<u32>(separate_addr) << desc.secondary_shift_left};
|
||||
return TexturePair(lhs_raw | rhs_raw, via_header_index);
|
||||
secondary = bind_cbuf(desc.secondary_cbuf_index, desc.count);
|
||||
has_secondary = true;
|
||||
secondary_offset = desc.secondary_cbuf_offset;
|
||||
primary_shift = desc.shift_left;
|
||||
secondary_shift = desc.secondary_shift_left;
|
||||
}
|
||||
}
|
||||
return TexturePair(gpu_memory.Read<u32>(addr), via_header_index);
|
||||
return [=](u32 index) {
|
||||
const u32 delta = index << desc.size_shift;
|
||||
const u32 raw = primary.Read(desc.cbuf_offset + delta);
|
||||
return TexturePair(has_secondary ? (raw << primary_shift) |
|
||||
(secondary.Read(secondary_offset + delta) << secondary_shift) : raw,
|
||||
via_header_index);
|
||||
};
|
||||
}};
|
||||
const auto add_image{[&](const auto& desc, bool blacklist) {
|
||||
const auto read_handle = prepare_handle(desc);
|
||||
for (u32 index = 0; index < desc.count; ++index) {
|
||||
const auto handle{read_handle(desc, index)};
|
||||
views.push_back({
|
||||
const auto handle{read_handle(index)};
|
||||
views[view_index++] = {
|
||||
.index = handle.first,
|
||||
.blacklist = blacklist,
|
||||
.id = {},
|
||||
});
|
||||
};
|
||||
}
|
||||
}};
|
||||
for (const auto& desc : info.texture_buffer_descriptors) {
|
||||
|
|
@ -161,17 +186,22 @@ void ComputePipeline::Configure(Tegra::Engines::KeplerCompute& kepler_compute,
|
|||
add_image(desc, false);
|
||||
}
|
||||
for (const auto& desc : info.texture_descriptors) {
|
||||
const auto read_handle = prepare_handle(desc);
|
||||
for (u32 index = 0; index < desc.count; ++index) {
|
||||
const auto handle{read_handle(desc, index)};
|
||||
views.push_back({handle.first});
|
||||
const auto handle{read_handle(index)};
|
||||
views[view_index++] = {handle.first};
|
||||
|
||||
VideoCommon::SamplerId sampler = texture_cache.GetComputeSamplerId(handle.second);
|
||||
samplers.push_back(sampler);
|
||||
const VideoCommon::SamplerId sampler = memoize_samplers
|
||||
? sampler_lookup->Get(handle.second, lookup_sampler)
|
||||
: lookup_sampler(handle.second);
|
||||
samplers[sampler_index++] = sampler;
|
||||
}
|
||||
}
|
||||
for (const auto& desc : info.image_descriptors) {
|
||||
add_image(desc, desc.is_written);
|
||||
}
|
||||
ASSERT(view_index == views.size());
|
||||
ASSERT(sampler_index == samplers.size());
|
||||
texture_cache.FillComputeImageViews(std::span(views.data(), views.size()));
|
||||
|
||||
buffer_cache.UnbindComputeTextureBuffers();
|
||||
|
|
|
|||
|
|
@ -5,11 +5,15 @@
|
|||
|
||||
#include <atomic>
|
||||
#include <condition_variable>
|
||||
#include <memory>
|
||||
#include <mutex>
|
||||
|
||||
#include <boost/container/small_vector.hpp>
|
||||
|
||||
#include "yuzu_common/common_types.h"
|
||||
#include "yuzu_common/thread_worker.h"
|
||||
#include "yuzu_shader_recompiler/shader_info.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/descriptor_lookup.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_buffer_cache.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_descriptor_pool.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_texture_cache.h"
|
||||
|
|
@ -55,6 +59,12 @@ private:
|
|||
vk::PipelineCache& pipeline_cache;
|
||||
GuestDescriptorQueue& guest_descriptor_queue;
|
||||
Shader::Info info;
|
||||
u32 num_textures{};
|
||||
size_t num_descriptor_entries{};
|
||||
|
||||
boost::container::small_vector<VideoCommon::ImageViewInOut, 64> descriptor_views;
|
||||
boost::container::small_vector<VideoCommon::SamplerId, 64> descriptor_samplers;
|
||||
std::unique_ptr<DrawDescriptorCache<VideoCommon::SamplerId>> sampler_lookup;
|
||||
|
||||
VideoCommon::ComputeUniformBufferSizes uniform_buffer_sizes{};
|
||||
|
||||
|
|
|
|||
|
|
@ -4,7 +4,6 @@
|
|||
#include <algorithm>
|
||||
#include <span>
|
||||
|
||||
#include <boost/container/small_vector.hpp>
|
||||
#include <boost/container/static_vector.hpp>
|
||||
|
||||
#include "yuzu_video_core/renderer_vulkan/pipeline_helper.h"
|
||||
|
|
@ -33,7 +32,6 @@
|
|||
|
||||
namespace Vulkan {
|
||||
namespace {
|
||||
using boost::container::small_vector;
|
||||
using boost::container::static_vector;
|
||||
using Shader::ImageBufferDescriptor;
|
||||
using Shader::Backend::SPIRV::RENDERAREA_LAYOUT_OFFSET;
|
||||
|
|
@ -45,7 +43,7 @@ using VideoCore::Surface::PixelFormatFromDepthFormat;
|
|||
using VideoCore::Surface::PixelFormatFromRenderTargetFormat;
|
||||
|
||||
constexpr size_t NUM_STAGES = Maxwell::MaxShaderStage;
|
||||
constexpr size_t MAX_IMAGE_ELEMENTS = 64;
|
||||
constexpr size_t INLINE_IMAGE_ELEMENTS = 64;
|
||||
|
||||
DescriptorLayoutBuilder MakeBuilder(const Device& device, std::span<const Shader::Info> infos,
|
||||
u32 num_color_attachments) {
|
||||
|
|
@ -92,14 +90,19 @@ DescriptorLayoutBuilder MakeBuilder(const Device& device, std::span<const Shader
|
|||
if (stages.at(index) == VK_SHADER_STAGE_FRAGMENT_BIT) {
|
||||
// Color attachments also count against maxPerStageResources.
|
||||
resources += num_color_attachments;
|
||||
const u64 storage_images =
|
||||
u64{Shader::NumDescriptors(infos[index].image_descriptors)} +
|
||||
Shader::NumDescriptors(infos[index].image_buffer_descriptors);
|
||||
can_widen = storage_images + num_storage_buffers + num_color_attachments <=
|
||||
device.DescriptorLimits().maxFragmentCombinedOutputResources;
|
||||
}
|
||||
can_widen = resources <= device.GetMaxPerStageResources();
|
||||
can_widen &= resources <= device.GetMaxPerStageResources();
|
||||
}
|
||||
if (!can_widen) {
|
||||
storage_buffer_stages = 0;
|
||||
}
|
||||
}
|
||||
DescriptorLayoutBuilder builder{device, storage_buffer_stages};
|
||||
DescriptorLayoutBuilder builder{device, storage_buffer_stages, num_color_attachments};
|
||||
for (size_t index = 0; index < infos.size(); ++index) {
|
||||
builder.Add(infos[index], stages.at(index));
|
||||
}
|
||||
|
|
@ -302,6 +305,7 @@ struct GraphicsPipeline::VtgState {
|
|||
Shader::Info info;
|
||||
/// Shader information of the resources bound by the guest
|
||||
Shader::Info guest_info;
|
||||
size_t num_descriptor_entries{};
|
||||
vk::ShaderModule module;
|
||||
vk::DescriptorSetLayout descriptor_set_layout;
|
||||
DescriptorAllocator descriptor_allocator;
|
||||
|
|
@ -370,6 +374,11 @@ GraphicsPipeline::GraphicsPipeline(
|
|||
enabled_uniform_buffer_masks[stage] = info->constant_buffer_mask;
|
||||
std::ranges::copy(info->constant_buffer_used_sizes, uniform_buffer_sizes[stage].begin());
|
||||
num_textures += Shader::NumDescriptors(info->texture_descriptors);
|
||||
num_descriptor_entries += NumDescriptorUpdateEntries(*info);
|
||||
num_image_elements += Shader::NumDescriptors(info->texture_buffer_descriptors);
|
||||
num_image_elements += Shader::NumDescriptors(info->image_buffer_descriptors);
|
||||
num_image_elements += Shader::NumDescriptors(info->texture_descriptors);
|
||||
num_image_elements += Shader::NumDescriptors(info->image_descriptors);
|
||||
}
|
||||
if (vtg_info) {
|
||||
static constexpr std::array<size_t, VtgPipelineInfo::NUM_COMPUTE_STAGES> compute_stages{
|
||||
|
|
@ -382,8 +391,15 @@ GraphicsPipeline::GraphicsPipeline(
|
|||
for (size_t index = 0; index < vtg->stages.size(); ++index) {
|
||||
VtgState::Stage& stage{vtg->stages[index]};
|
||||
stage.info = vtg_info->infos[index];
|
||||
stage.num_descriptor_entries = NumDescriptorUpdateEntries(stage.info);
|
||||
stage.guest_info = stage.info;
|
||||
stage.guest_info.storage_buffers_descriptors.resize(guest_storage_buffers[index]);
|
||||
// The compute stages retain guest images absent from the generated vertex shader.
|
||||
num_textures += Shader::NumDescriptors(stage.guest_info.texture_descriptors);
|
||||
num_image_elements += Shader::NumDescriptors(stage.guest_info.texture_buffer_descriptors);
|
||||
num_image_elements += Shader::NumDescriptors(stage.guest_info.image_buffer_descriptors);
|
||||
num_image_elements += Shader::NumDescriptors(stage.guest_info.texture_descriptors);
|
||||
num_image_elements += Shader::NumDescriptors(stage.guest_info.image_descriptors);
|
||||
stage.module = std::move(vtg_info->modules[index]);
|
||||
vtg->uniform_masks[compute_stages[index]] = stage.info.constant_buffer_mask;
|
||||
std::ranges::copy(stage.info.constant_buffer_used_sizes,
|
||||
|
|
@ -397,6 +413,8 @@ GraphicsPipeline::GraphicsPipeline(
|
|||
vtg->scratch = vtg_info->scratch;
|
||||
vtg->staging_pool = vtg_info->staging_pool;
|
||||
}
|
||||
descriptor_views.resize(num_image_elements);
|
||||
descriptor_samplers.resize(num_textures);
|
||||
// Number of leading components the vertex shader may read from each attribute. Used to keep
|
||||
// vertex formats within their binding stride without discarding data the shader reads.
|
||||
const Shader::Info& vertex_info{stage_infos[0]};
|
||||
|
|
@ -505,13 +523,24 @@ void GraphicsPipeline::AddTransition(GraphicsPipeline* transition,
|
|||
|
||||
template <typename Spec>
|
||||
void GraphicsPipeline::ConfigureImpl(bool is_indexed) {
|
||||
std::array<VideoCommon::ImageViewInOut, MAX_IMAGE_ELEMENTS> views;
|
||||
std::array<VideoCommon::SamplerId, MAX_IMAGE_ELEMENTS> samplers;
|
||||
size_t sampler_index{};
|
||||
auto& views = descriptor_views;
|
||||
auto& samplers = descriptor_samplers;
|
||||
size_t view_index{};
|
||||
size_t sampler_index{};
|
||||
|
||||
texture_cache.SynchronizeGraphicsDescriptors();
|
||||
|
||||
if (num_textures > INLINE_IMAGE_ELEMENTS) {
|
||||
if (!sampler_lookup) {
|
||||
sampler_lookup = std::make_unique<DrawDescriptorCache<VideoCommon::SamplerId>>();
|
||||
}
|
||||
sampler_lookup->Reset();
|
||||
}
|
||||
const bool memoize_samplers = sampler_lookup && sampler_lookup->Enabled();
|
||||
const auto lookup_sampler = [&](u32 index) {
|
||||
return texture_cache.GetGraphicsSamplerId(index);
|
||||
};
|
||||
|
||||
buffer_cache.SetUniformBuffersState(enabled_uniform_buffer_masks, &uniform_buffer_sizes);
|
||||
|
||||
const auto& regs{maxwell3d->regs};
|
||||
|
|
@ -529,30 +558,38 @@ void GraphicsPipeline::ConfigureImpl(bool is_indexed) {
|
|||
}
|
||||
}
|
||||
const auto& cbufs{maxwell3d->state.shader_stages[stage].const_buffers};
|
||||
const auto read_handle{[&](const auto& desc, u32 index) {
|
||||
DescriptorCbufReader<18> cbuf_reader;
|
||||
const auto bind_cbuf = [&](u32 bank, u32 count) {
|
||||
return cbuf_reader.Bind(*gpu_memory, bank, cbufs[bank].address, cbufs[bank].size, count);
|
||||
};
|
||||
const auto prepare_handle{[&](const auto& desc) {
|
||||
ASSERT(cbufs[desc.cbuf_index].enabled);
|
||||
const u32 index_offset{index << desc.size_shift};
|
||||
const u32 offset{desc.cbuf_offset + index_offset};
|
||||
const GPUVAddr addr{cbufs[desc.cbuf_index].address + offset};
|
||||
if constexpr (std::is_same_v<decltype(desc), const Shader::TextureDescriptor&> ||
|
||||
std::is_same_v<decltype(desc), const Shader::TextureBufferDescriptor&>) {
|
||||
const auto primary = bind_cbuf(desc.cbuf_index, desc.count);
|
||||
auto secondary = primary;
|
||||
bool has_secondary{};
|
||||
u32 secondary_offset{}, primary_shift{}, secondary_shift{};
|
||||
if constexpr (requires { desc.has_secondary; }) {
|
||||
if (desc.has_secondary) {
|
||||
ASSERT(cbufs[desc.secondary_cbuf_index].enabled);
|
||||
const u32 second_offset{desc.secondary_cbuf_offset + index_offset};
|
||||
const GPUVAddr separate_addr{cbufs[desc.secondary_cbuf_index].address +
|
||||
second_offset};
|
||||
const u32 lhs_raw{gpu_memory->Read<u32>(addr) << desc.shift_left};
|
||||
const u32 rhs_raw{gpu_memory->Read<u32>(separate_addr)
|
||||
<< desc.secondary_shift_left};
|
||||
const u32 raw{lhs_raw | rhs_raw};
|
||||
return TexturePair(raw, via_header_index);
|
||||
secondary = bind_cbuf(desc.secondary_cbuf_index, desc.count);
|
||||
has_secondary = true;
|
||||
secondary_offset = desc.secondary_cbuf_offset;
|
||||
primary_shift = desc.shift_left;
|
||||
secondary_shift = desc.secondary_shift_left;
|
||||
}
|
||||
}
|
||||
return TexturePair(gpu_memory->Read<u32>(addr), via_header_index);
|
||||
return [=](u32 index) {
|
||||
const u32 delta = index << desc.size_shift;
|
||||
const u32 raw = primary.Read(desc.cbuf_offset + delta);
|
||||
return TexturePair(has_secondary ? (raw << primary_shift) |
|
||||
(secondary.Read(secondary_offset + delta) << secondary_shift) : raw,
|
||||
via_header_index);
|
||||
};
|
||||
}};
|
||||
const auto add_image{[&](const auto& desc, bool blacklist) LAMBDA_FORCEINLINE {
|
||||
const auto read_handle = prepare_handle(desc);
|
||||
for (u32 index = 0; index < desc.count; ++index) {
|
||||
const auto handle{read_handle(desc, index)};
|
||||
const auto handle{read_handle(index)};
|
||||
views[view_index++] = {
|
||||
.index = handle.first,
|
||||
.blacklist = blacklist,
|
||||
|
|
@ -571,11 +608,14 @@ void GraphicsPipeline::ConfigureImpl(bool is_indexed) {
|
|||
}
|
||||
}
|
||||
for (const auto& desc : info.texture_descriptors) {
|
||||
const auto read_handle = prepare_handle(desc);
|
||||
for (u32 index = 0; index < desc.count; ++index) {
|
||||
const auto handle{read_handle(desc, index)};
|
||||
const auto handle{read_handle(index)};
|
||||
views[view_index++] = {handle.first};
|
||||
|
||||
VideoCommon::SamplerId sampler{texture_cache.GetGraphicsSamplerId(handle.second)};
|
||||
const VideoCommon::SamplerId sampler = memoize_samplers
|
||||
? sampler_lookup->Get(handle.second, lookup_sampler)
|
||||
: lookup_sampler(handle.second);
|
||||
samplers[sampler_index++] = sampler;
|
||||
}
|
||||
}
|
||||
|
|
@ -600,7 +640,9 @@ void GraphicsPipeline::ConfigureImpl(bool is_indexed) {
|
|||
if constexpr (Spec::enabled_stages[4]) {
|
||||
config_stage(4);
|
||||
}
|
||||
texture_cache.FillGraphicsImageViews<Spec::has_images>(std::span(views.data(), view_index));
|
||||
ASSERT(view_index == views.size());
|
||||
ASSERT(sampler_index == samplers.size());
|
||||
texture_cache.FillGraphicsImageViews<Spec::has_images>(std::span(views.data(), views.size()));
|
||||
|
||||
VideoCommon::ImageViewInOut* texture_buffer_it{views.data()};
|
||||
const auto bind_stage_info{[&](size_t stage) LAMBDA_FORCEINLINE {
|
||||
|
|
@ -657,7 +699,7 @@ void GraphicsPipeline::ConfigureImpl(bool is_indexed) {
|
|||
buffer_cache.UpdateGraphicsBuffers(is_indexed);
|
||||
buffer_cache.BindHostGeometryBuffers(is_indexed);
|
||||
|
||||
guest_descriptor_queue.Acquire();
|
||||
guest_descriptor_queue.Acquire(num_descriptor_entries);
|
||||
|
||||
RescalingPushConstant rescaling;
|
||||
RenderAreaPushConstant render_area;
|
||||
|
|
@ -704,8 +746,8 @@ void GraphicsPipeline::ConfigureVtg(bool is_indexed) {
|
|||
VtgState& state{*vtg};
|
||||
state.index_count = 0;
|
||||
|
||||
std::array<VideoCommon::ImageViewInOut, MAX_IMAGE_ELEMENTS> views;
|
||||
std::array<VideoCommon::SamplerId, MAX_IMAGE_ELEMENTS> samplers;
|
||||
auto& views = descriptor_views;
|
||||
auto& samplers = descriptor_samplers;
|
||||
size_t sampler_index{};
|
||||
size_t view_index{};
|
||||
|
||||
|
|
@ -779,7 +821,9 @@ void GraphicsPipeline::ConfigureVtg(bool is_indexed) {
|
|||
add_image(desc, desc.is_written);
|
||||
}
|
||||
}
|
||||
texture_cache.FillGraphicsImageViews<true>(std::span(views.data(), view_index));
|
||||
ASSERT(view_index == views.size());
|
||||
ASSERT(sampler_index == samplers.size());
|
||||
texture_cache.FillGraphicsImageViews<true>(std::span(views.data(), views.size()));
|
||||
|
||||
VideoCommon::ImageViewInOut* texture_buffer_it{views.data()};
|
||||
for (const auto& [stage, info_pointer] : stages) {
|
||||
|
|
@ -932,7 +976,7 @@ void GraphicsPipeline::ConfigureVtg(bool is_indexed) {
|
|||
const auto run_stage{[&](size_t index, size_t guest_stage,
|
||||
std::span<const VtgState::Buffer> buffers, u32 invocations) {
|
||||
VtgState::Stage& stage{state.stages[index]};
|
||||
guest_descriptor_queue.Acquire();
|
||||
guest_descriptor_queue.Acquire(stage.num_descriptor_entries);
|
||||
buffer_cache.BindHostStageUniformAndStorageBuffers(guest_stage);
|
||||
for (const VtgState::Buffer& buffer : buffers) {
|
||||
guest_descriptor_queue.AddBuffer(buffer.buffer, buffer.offset, buffer.size);
|
||||
|
|
@ -1010,7 +1054,7 @@ void GraphicsPipeline::ConfigureVtg(bool is_indexed) {
|
|||
run_stage(VtgPipelineInfo::GEOMETRY, 3, geometry_buffers, geometry_invocations);
|
||||
|
||||
// The generated vertex shader reads the vertices written by the geometry stage
|
||||
guest_descriptor_queue.Acquire();
|
||||
guest_descriptor_queue.Acquire(num_descriptor_entries);
|
||||
guest_descriptor_queue.AddBuffer(geometry_buffer, geometry_offset, geometry_size);
|
||||
|
||||
RescalingPushConstant rescaling;
|
||||
|
|
@ -1551,7 +1595,15 @@ void GraphicsPipeline::Validate() {
|
|||
num_images += Shader::NumDescriptors(info.texture_descriptors);
|
||||
num_images += Shader::NumDescriptors(info.image_descriptors);
|
||||
}
|
||||
ASSERT(num_images <= MAX_IMAGE_ELEMENTS);
|
||||
if (vtg) {
|
||||
for (const auto& stage : vtg->stages) {
|
||||
num_images += Shader::NumDescriptors(stage.guest_info.texture_buffer_descriptors);
|
||||
num_images += Shader::NumDescriptors(stage.guest_info.image_buffer_descriptors);
|
||||
num_images += Shader::NumDescriptors(stage.guest_info.texture_descriptors);
|
||||
num_images += Shader::NumDescriptors(stage.guest_info.image_descriptors);
|
||||
}
|
||||
}
|
||||
ASSERT(num_images == num_image_elements);
|
||||
}
|
||||
|
||||
} // namespace Vulkan
|
||||
|
|
|
|||
|
|
@ -13,11 +13,14 @@
|
|||
#include <type_traits>
|
||||
#include <utility>
|
||||
|
||||
#include <boost/container/small_vector.hpp>
|
||||
|
||||
#include "yuzu_common/thread_worker.h"
|
||||
#include "yuzu_shader_recompiler/shader_info.h"
|
||||
#include "yuzu_shader_recompiler/vtg_as_compute.h"
|
||||
#include "yuzu_video_core/engines/maxwell_3d.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/fixed_pipeline_state.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/descriptor_lookup.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_buffer_cache.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_descriptor_pool.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_texture_cache.h"
|
||||
|
|
@ -31,6 +34,7 @@ namespace Vulkan {
|
|||
|
||||
struct GraphicsPipelineCacheKey {
|
||||
std::array<u64, 6> unique_hashes;
|
||||
std::array<std::array<u32, 18>, 5> cbuf_sizes;
|
||||
FixedPipelineState state;
|
||||
|
||||
size_t Hash() const noexcept;
|
||||
|
|
@ -42,7 +46,7 @@ struct GraphicsPipelineCacheKey {
|
|||
}
|
||||
|
||||
size_t Size() const noexcept {
|
||||
return sizeof(unique_hashes) + state.Size();
|
||||
return sizeof(unique_hashes) + sizeof(cbuf_sizes) + state.Size();
|
||||
}
|
||||
};
|
||||
static_assert(std::has_unique_object_representations_v<GraphicsPipelineCacheKey>);
|
||||
|
|
@ -219,6 +223,13 @@ private:
|
|||
std::array<u32, 5> enabled_uniform_buffer_masks{};
|
||||
VideoCommon::UniformBufferSizes uniform_buffer_sizes{};
|
||||
u32 num_textures{};
|
||||
u32 num_image_elements{};
|
||||
size_t num_descriptor_entries{};
|
||||
|
||||
// Used only by the GPU thread; queued commands use GuestDescriptorQueue's own copies.
|
||||
boost::container::small_vector<VideoCommon::ImageViewInOut, 64> descriptor_views;
|
||||
boost::container::small_vector<VideoCommon::SamplerId, 64> descriptor_samplers;
|
||||
std::unique_ptr<DrawDescriptorCache<VideoCommon::SamplerId>> sampler_lookup;
|
||||
|
||||
vk::DescriptorSetLayout descriptor_set_layout;
|
||||
DescriptorAllocator descriptor_allocator;
|
||||
|
|
|
|||
|
|
@ -6,6 +6,8 @@
|
|||
#include <fstream>
|
||||
#include <memory>
|
||||
#include <thread>
|
||||
#include <unordered_set>
|
||||
#include "yuzu_common/scope_exit.h"
|
||||
#include <vector>
|
||||
|
||||
#include "yuzu_common/bit_cast.h"
|
||||
|
|
@ -54,7 +56,7 @@ using VideoCommon::FileEnvironment;
|
|||
using VideoCommon::GenericEnvironment;
|
||||
using VideoCommon::GraphicsEnvironment;
|
||||
|
||||
constexpr u32 CACHE_VERSION = 11;
|
||||
constexpr u32 CACHE_VERSION = 14;
|
||||
constexpr std::array<char, 8> VULKAN_CACHE_MAGIC_NUMBER{'y', 'u', 'z', 'u', 'v', 'k', 'c', 'h'};
|
||||
|
||||
template <typename Container>
|
||||
|
|
@ -471,6 +473,8 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
|
|||
.supported_spirv = device.SupportedSpirvVersion(),
|
||||
.unified_descriptor_binding = true,
|
||||
.support_descriptor_aliasing = device.IsDescriptorAliasingSupported(),
|
||||
.support_sampled_image_array_nonuniform_indexing =
|
||||
device.IsSampledImageArrayNonUniformIndexingSupported(),
|
||||
.disable_uniform_buffer_descriptor_aliasing =
|
||||
device.IsDescriptorAliasingSupported() && device.GetDriverID() == VK_DRIVER_ID_MOLTENVK,
|
||||
.support_int8 = device.IsInt8Supported(),
|
||||
|
|
@ -548,6 +552,7 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
|
|||
.min_ssbo_alignment = static_cast<u32>(device.GetStorageBufferAlignment()),
|
||||
.support_geometry_shader_passthrough = device.IsNvGeometryShaderPassthroughSupported(),
|
||||
.support_conditional_barrier = device.SupportsConditionalBarriers(),
|
||||
.max_bindless_descriptors_per_stage = device.MaxBindlessDescriptorsPerStage(),
|
||||
};
|
||||
|
||||
if (device.GetMaxVertexInputAttributes() < Maxwell::NumVertexAttributes) {
|
||||
|
|
@ -584,6 +589,11 @@ GraphicsPipeline* PipelineCache::CurrentGraphicsPipeline() {
|
|||
return nullptr;
|
||||
}
|
||||
graphics_key.state.Refresh(*maxwell3d, dynamic_features);
|
||||
graphics_cbuf_size_state.Refresh(graphics_key.cbuf_sizes, graphics_key.unique_hashes,
|
||||
cbuf_size_dependencies, [&](size_t stage, size_t bank) {
|
||||
const auto& buffer = maxwell3d->state.shader_stages[stage].const_buffers[bank];
|
||||
return buffer.enabled ? buffer.size : 0;
|
||||
});
|
||||
|
||||
if (current_pipeline) {
|
||||
GraphicsPipeline* const next{current_pipeline->Next(graphics_key)};
|
||||
|
|
@ -601,11 +611,17 @@ ComputePipeline* PipelineCache::CurrentComputePipeline() {
|
|||
return nullptr;
|
||||
}
|
||||
const auto& qmd{kepler_compute->launch_description};
|
||||
const ComputePipelineCacheKey key{
|
||||
ComputePipelineCacheKey key{
|
||||
.unique_hash = shader->unique_hash,
|
||||
.shared_memory_size = qmd.shared_alloc,
|
||||
.workgroup_size{qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z},
|
||||
};
|
||||
const u32 size_mask = qmd.const_buffer_enable_mask.Value() &
|
||||
cbuf_size_dependencies.Mask(shader->unique_hash);
|
||||
for (size_t bank = 0; bank < key.cbuf_sizes.size(); ++bank) {
|
||||
key.cbuf_sizes[bank] = ((size_mask >> bank) & 1) != 0
|
||||
? qmd.const_buffer_config[bank].size.Value() : 0;
|
||||
}
|
||||
const auto [pair, is_new]{compute_cache.try_emplace(key)};
|
||||
auto& pipeline{pair->second};
|
||||
if (!is_new) {
|
||||
|
|
@ -641,14 +657,29 @@ void PipelineCache::LoadDiskResources(u64 title_id, std::stop_token stop_loading
|
|||
bool has_loaded{};
|
||||
std::unique_ptr<PipelineStatistics> statistics;
|
||||
} state;
|
||||
// Cache parsing/allocation can throw after work capturing this stack state is queued.
|
||||
// Drain active work before state and the progress callback can leave scope.
|
||||
SCOPE_EXIT { workers.WaitForRequests(); };
|
||||
|
||||
if (device.IsKhrPipelineExecutablePropertiesEnabled()) {
|
||||
state.statistics = std::make_unique<PipelineStatistics>(device);
|
||||
}
|
||||
|
||||
// Earlier records included sizes for every bound bank, even when the shader never
|
||||
// queried them. Normalize and deduplicate them without discarding the existing shader cache.
|
||||
std::unordered_set<ComputePipelineCacheKey> loaded_compute_keys;
|
||||
std::unordered_set<GraphicsPipelineCacheKey> loaded_graphics_keys;
|
||||
const auto load_compute{[&](std::ifstream& file, FileEnvironment env) {
|
||||
ComputePipelineCacheKey key;
|
||||
file.read(reinterpret_cast<char*>(&key), sizeof(key));
|
||||
|
||||
cbuf_size_dependencies.Observe(key.unique_hash, env.CbufSizeMask());
|
||||
CbufSizeDependencies::Normalize(key.cbuf_sizes,
|
||||
cbuf_size_dependencies.Mask(key.unique_hash));
|
||||
if (!loaded_compute_keys.insert(key).second) {
|
||||
return;
|
||||
}
|
||||
|
||||
workers.QueueWork([this, key, env_ = std::move(env), &state, &callback]() mutable {
|
||||
ShaderPools pools;
|
||||
auto pipeline{CreateComputePipeline(pools, key, env_, state.statistics.get(), false)};
|
||||
|
|
@ -667,7 +698,7 @@ void PipelineCache::LoadDiskResources(u64 title_id, std::stop_token stop_loading
|
|||
std::vector<FileEnvironment> envs) {
|
||||
workers.QueueWork([this, key, envs_ = std::move(envs), &state, &callback]() mutable {
|
||||
ShaderPools pools;
|
||||
boost::container::static_vector<Shader::Environment*, 5> env_ptrs;
|
||||
boost::container::static_vector<Shader::Environment*, Maxwell::MaxShaderProgram> env_ptrs;
|
||||
for (auto& env : envs_) {
|
||||
env_ptrs.push_back(&env);
|
||||
}
|
||||
|
|
@ -697,6 +728,29 @@ void PipelineCache::LoadDiskResources(u64 title_id, std::stop_token stop_loading
|
|||
GraphicsPipelineCacheKey key;
|
||||
file.read(reinterpret_cast<char*>(&key), sizeof(key));
|
||||
|
||||
size_t env_index{};
|
||||
for (u64 hash : key.unique_hashes) {
|
||||
if (hash != 0) {
|
||||
if (env_index >= envs.size()) {
|
||||
LOG_ERROR(Render_Vulkan, "Cached graphics shader environment count mismatch");
|
||||
return;
|
||||
}
|
||||
cbuf_size_dependencies.Observe(hash, envs[env_index++].CbufSizeMask());
|
||||
}
|
||||
}
|
||||
if (env_index != envs.size()) {
|
||||
LOG_ERROR(Render_Vulkan, "Cached graphics shader environment count mismatch");
|
||||
return;
|
||||
}
|
||||
for (size_t stage = 0; stage < key.cbuf_sizes.size(); ++stage) {
|
||||
const u32 mask = cbuf_size_dependencies.Mask(key.unique_hashes[stage + 1]) |
|
||||
(stage == 0 ? cbuf_size_dependencies.Mask(key.unique_hashes[0]) : 0);
|
||||
CbufSizeDependencies::Normalize(key.cbuf_sizes[stage], mask);
|
||||
}
|
||||
if (!loaded_graphics_keys.insert(key).second) {
|
||||
return;
|
||||
}
|
||||
|
||||
if ((key.state.extended_dynamic_state != 0) !=
|
||||
dynamic_features.has_extended_dynamic_state ||
|
||||
(key.state.extended_dynamic_state_2 != 0) !=
|
||||
|
|
@ -1061,6 +1115,10 @@ std::shared_ptr<GraphicsPipeline> PipelineCache::CreateGraphicsPipeline() {
|
|||
std::shared_ptr<GraphicsPipeline> pipeline{
|
||||
CreateGraphicsPipeline(main_pools, graphics_key, environments.Span(), nullptr, true,
|
||||
is_hle && twin ? &skip_build : nullptr)};
|
||||
for (size_t index = 0; index < graphics_key.unique_hashes.size(); ++index) {
|
||||
cbuf_size_dependencies.Observe(graphics_key.unique_hashes[index],
|
||||
environments.envs[index].CbufSizeMask());
|
||||
}
|
||||
if (reuse_twin) {
|
||||
twin->SetSharedKey(graphics_key);
|
||||
pipeline = twin;
|
||||
|
|
@ -1092,6 +1150,7 @@ std::unique_ptr<ComputePipeline> PipelineCache::CreateComputePipeline(
|
|||
|
||||
main_pools.ReleaseContents();
|
||||
auto pipeline{CreateComputePipeline(main_pools, key, env, nullptr, true)};
|
||||
cbuf_size_dependencies.Observe(key.unique_hash, env.CbufSizeMask());
|
||||
if (!pipeline || pipeline_cache_filename.empty()) {
|
||||
return pipeline;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -24,6 +24,7 @@
|
|||
#include "yuzu_video_core/engines/maxwell_3d.h"
|
||||
#include "yuzu_video_core/host1x/gpu_device_memory_manager.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/fixed_pipeline_state.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/cbuf_size_dependencies.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_buffer_cache.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_compute_pipeline.h"
|
||||
#include "yuzu_video_core/renderer_vulkan/vk_graphics_pipeline.h"
|
||||
|
|
@ -50,6 +51,7 @@ struct ComputePipelineCacheKey {
|
|||
u64 unique_hash;
|
||||
u32 shared_memory_size;
|
||||
std::array<u32, 3> workgroup_size;
|
||||
std::array<u32, 8> cbuf_sizes;
|
||||
|
||||
size_t Hash() const noexcept;
|
||||
|
||||
|
|
@ -159,6 +161,8 @@ private:
|
|||
|
||||
GraphicsPipelineCacheKey graphics_key{};
|
||||
GraphicsPipeline* current_pipeline{};
|
||||
CbufSizeDependencies cbuf_size_dependencies;
|
||||
GraphicsCbufSizeState<5, 18> graphics_cbuf_size_state;
|
||||
|
||||
std::unordered_map<ComputePipelineCacheKey, std::unique_ptr<ComputePipeline>> compute_cache;
|
||||
// Shared: an HLE macro key without replacements maps to the pipeline of its non-HLE twin.
|
||||
|
|
|
|||
|
|
@ -13,32 +13,19 @@
|
|||
namespace Vulkan {
|
||||
|
||||
UpdateDescriptorQueue::UpdateDescriptorQueue(const Device& device_, Scheduler& scheduler_)
|
||||
: device{device_}, scheduler{scheduler_} {
|
||||
payload_start = payload.data();
|
||||
payload_cursor = payload.data();
|
||||
}
|
||||
: device{device_}, scheduler{scheduler_} {}
|
||||
|
||||
UpdateDescriptorQueue::~UpdateDescriptorQueue() = default;
|
||||
|
||||
void UpdateDescriptorQueue::TickFrame() {
|
||||
if (++frame_index >= FRAMES_IN_FLIGHT) {
|
||||
frame_index = 0;
|
||||
}
|
||||
payload_start = payload.data() + frame_index * FRAME_PAYLOAD_SIZE;
|
||||
payload_cursor = payload_start;
|
||||
payload.TickFrame();
|
||||
}
|
||||
|
||||
void UpdateDescriptorQueue::Acquire() {
|
||||
// Minimum number of entries required.
|
||||
// This is the maximum number of entries a single draw call might use.
|
||||
static constexpr size_t MIN_ENTRIES = 0x400;
|
||||
|
||||
if (std::distance(payload_start, payload_cursor) + MIN_ENTRIES >= FRAME_PAYLOAD_SIZE) {
|
||||
LOG_WARNING(Render_Vulkan, "Payload overflow, waiting for worker thread");
|
||||
void UpdateDescriptorQueue::Acquire(size_t num_entries) {
|
||||
payload.Acquire(num_entries, [&] {
|
||||
LOG_DEBUG(Render_Vulkan, "Recycling descriptor payload after waiting for worker");
|
||||
scheduler.WaitWorker();
|
||||
payload_cursor = payload_start;
|
||||
}
|
||||
upload_start = payload_cursor;
|
||||
});
|
||||
}
|
||||
|
||||
} // namespace Vulkan
|
||||
|
|
|
|||
|
|
@ -3,7 +3,7 @@
|
|||
|
||||
#pragma once
|
||||
|
||||
#include <array>
|
||||
#include "yuzu_video_core/renderer_vulkan/descriptor_payload.h"
|
||||
|
||||
#include "yuzu_video_core/vulkan_common/vulkan_wrapper.h"
|
||||
|
||||
|
|
@ -33,7 +33,6 @@ class UpdateDescriptorQueue final {
|
|||
// provide up to 3 swapchain images.
|
||||
static constexpr size_t FRAMES_IN_FLIGHT = 8;
|
||||
static constexpr size_t FRAME_PAYLOAD_SIZE = 0x20000;
|
||||
static constexpr size_t PAYLOAD_SIZE = FRAME_PAYLOAD_SIZE * FRAMES_IN_FLIGHT;
|
||||
|
||||
public:
|
||||
explicit UpdateDescriptorQueue(const Device& device_, Scheduler& scheduler_);
|
||||
|
|
@ -41,49 +40,46 @@ public:
|
|||
|
||||
void TickFrame();
|
||||
|
||||
void Acquire();
|
||||
void Acquire(size_t num_entries);
|
||||
|
||||
const DescriptorUpdateEntry* UpdateData() const noexcept {
|
||||
return upload_start;
|
||||
return payload.Data();
|
||||
}
|
||||
|
||||
void AddSampledImage(VkImageView image_view, VkSampler sampler) {
|
||||
*(payload_cursor++) = VkDescriptorImageInfo{
|
||||
payload.Append(VkDescriptorImageInfo{
|
||||
.sampler = sampler,
|
||||
.imageView = image_view,
|
||||
.imageLayout = VK_IMAGE_LAYOUT_GENERAL,
|
||||
};
|
||||
});
|
||||
}
|
||||
|
||||
void AddImage(VkImageView image_view) {
|
||||
*(payload_cursor++) = VkDescriptorImageInfo{
|
||||
payload.Append(VkDescriptorImageInfo{
|
||||
.sampler = VK_NULL_HANDLE,
|
||||
.imageView = image_view,
|
||||
.imageLayout = VK_IMAGE_LAYOUT_GENERAL,
|
||||
};
|
||||
});
|
||||
}
|
||||
|
||||
void AddBuffer(VkBuffer buffer, VkDeviceSize offset, VkDeviceSize size) {
|
||||
*(payload_cursor++) = VkDescriptorBufferInfo{
|
||||
payload.Append(VkDescriptorBufferInfo{
|
||||
.buffer = buffer,
|
||||
.offset = offset,
|
||||
.range = size,
|
||||
};
|
||||
});
|
||||
}
|
||||
|
||||
void AddTexelBuffer(VkBufferView texel_buffer) {
|
||||
*(payload_cursor++) = texel_buffer;
|
||||
payload.Append(texel_buffer);
|
||||
}
|
||||
|
||||
private:
|
||||
const Device& device;
|
||||
Scheduler& scheduler;
|
||||
|
||||
size_t frame_index{0};
|
||||
DescriptorUpdateEntry* payload_cursor = nullptr;
|
||||
DescriptorUpdateEntry* payload_start = nullptr;
|
||||
const DescriptorUpdateEntry* upload_start = nullptr;
|
||||
std::array<DescriptorUpdateEntry, PAYLOAD_SIZE> payload;
|
||||
DescriptorPayload<DescriptorUpdateEntry, FRAME_PAYLOAD_SIZE, FRAMES_IN_FLIGHT> payload;
|
||||
|
||||
};
|
||||
|
||||
// TODO: should these be separate classes instead?
|
||||
|
|
|
|||
|
|
@ -35,6 +35,25 @@ static u64 MakeCbufKey(u32 index, u32 offset) {
|
|||
return (static_cast<u64>(index) << 32) | offset;
|
||||
}
|
||||
|
||||
static u32 CbufSizeQueryMask(const std::unordered_map<u64, u32>& values) {
|
||||
u32 mask{};
|
||||
for (const auto& [key, value] : values) {
|
||||
const u32 bank = static_cast<u32>(key >> 32);
|
||||
if (static_cast<u32>(key) == UINT32_MAX && bank < 32) {
|
||||
mask |= 1U << bank;
|
||||
}
|
||||
}
|
||||
return mask;
|
||||
}
|
||||
|
||||
u32 GenericEnvironment::CbufSizeMask() const noexcept {
|
||||
return CbufSizeQueryMask(cbuf_values);
|
||||
}
|
||||
|
||||
u32 FileEnvironment::CbufSizeMask() const noexcept {
|
||||
return CbufSizeQueryMask(cbuf_values);
|
||||
}
|
||||
|
||||
static Shader::TextureType ConvertTextureType(
|
||||
const std::optional<Tegra::Texture::TICEntry>& optional_entry) {
|
||||
if (!optional_entry) {
|
||||
|
|
@ -341,6 +360,14 @@ GraphicsEnvironment::GraphicsEnvironment(Tegra::Engines::Maxwell3D& maxwell3d_,
|
|||
maxwell3d->engine_state == Tegra::Engines::Maxwell3D::EngineHint::OnHLEMacro;
|
||||
}
|
||||
|
||||
u32 GraphicsEnvironment::ReadCbufSize(u32 cbuf_index) {
|
||||
const auto& cbuf{maxwell3d->state.shader_stages[stage_index].const_buffers[cbuf_index]};
|
||||
const u32 size = cbuf.enabled ? cbuf.size : 0;
|
||||
// Reserved offset outside the guest address space; keeps cache layout unchanged.
|
||||
cbuf_values.emplace(MakeCbufKey(cbuf_index, UINT32_MAX), size);
|
||||
return size;
|
||||
}
|
||||
|
||||
u32 GraphicsEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
|
||||
const auto& cbuf{maxwell3d->state.shader_stages[stage_index].const_buffers[cbuf_index]};
|
||||
ASSERT(cbuf.enabled);
|
||||
|
|
@ -423,6 +450,14 @@ ComputeEnvironment::ComputeEnvironment(Tegra::Engines::KeplerCompute& kepler_com
|
|||
workgroup_size = {qmd.block_dim_x, qmd.block_dim_y, qmd.block_dim_z};
|
||||
}
|
||||
|
||||
u32 ComputeEnvironment::ReadCbufSize(u32 cbuf_index) {
|
||||
const auto& qmd{kepler_compute->launch_description};
|
||||
const bool enabled = ((qmd.const_buffer_enable_mask.Value() >> cbuf_index) & 1) != 0;
|
||||
const u32 size = enabled ? qmd.const_buffer_config[cbuf_index].size.Value() : 0;
|
||||
cbuf_values.emplace(MakeCbufKey(cbuf_index, UINT32_MAX), size);
|
||||
return size;
|
||||
}
|
||||
|
||||
u32 ComputeEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
|
||||
const auto& qmd{kepler_compute->launch_description};
|
||||
ASSERT(((qmd.const_buffer_enable_mask.Value() >> cbuf_index) & 1) != 0);
|
||||
|
|
@ -482,6 +517,33 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
|
|||
.read(reinterpret_cast<char*>(&read_highest), sizeof(read_highest))
|
||||
.read(reinterpret_cast<char*>(&viewport_transform_state), sizeof(viewport_transform_state))
|
||||
.read(reinterpret_cast<char*>(&stage), sizeof(stage));
|
||||
// A crash can leave the last cache entry incomplete. Validate serialized sizes before
|
||||
// allocating memory; otherwise arbitrary bytes can turn into a huge allocation request.
|
||||
if (!file || stage > Shader::Stage::VertexA || code_size % sizeof(u64) != 0) {
|
||||
throw std::ios_base::failure("Invalid shader cache environment header");
|
||||
}
|
||||
const auto position = file.tellg();
|
||||
file.seekg(0, std::ios::end);
|
||||
const auto end = file.tellg();
|
||||
file.seekg(position);
|
||||
if (!file || position < 0 || end < position) {
|
||||
throw std::ios_base::failure("Invalid shader cache file bounds");
|
||||
}
|
||||
u64 remaining = static_cast<u64>(end - position);
|
||||
if (code_size > remaining) {
|
||||
throw std::ios_base::failure("Invalid shader cache code size");
|
||||
}
|
||||
remaining -= code_size;
|
||||
const auto consume_entries = [&remaining](u64 count, size_t entry_size) {
|
||||
if (count > remaining / entry_size) {
|
||||
throw std::ios_base::failure("Invalid shader cache entry count");
|
||||
}
|
||||
remaining -= count * entry_size;
|
||||
};
|
||||
consume_entries(num_texture_types, sizeof(u32) + sizeof(Shader::TextureType));
|
||||
consume_entries(num_texture_pixel_formats, sizeof(u32) + sizeof(Shader::TexturePixelFormat));
|
||||
consume_entries(num_cbuf_values, sizeof(u64) + sizeof(u32));
|
||||
consume_entries(num_cbuf_replacement_values, sizeof(u64) + sizeof(Shader::ReplaceConstant));
|
||||
code.resize(Common::DivCeil(code_size, sizeof(u64)));
|
||||
file.read(reinterpret_cast<char*>(code.data()), code_size);
|
||||
for (size_t i = 0; i < num_texture_types; ++i) {
|
||||
|
|
@ -523,6 +585,9 @@ void FileEnvironment::Deserialize(std::ifstream& file) {
|
|||
file.read(reinterpret_cast<char*>(&gp_passthrough_mask), sizeof(gp_passthrough_mask));
|
||||
}
|
||||
}
|
||||
if (!file) {
|
||||
throw std::ios_base::failure("Truncated shader cache environment");
|
||||
}
|
||||
is_proprietary_driver = texture_bound == 2;
|
||||
}
|
||||
|
||||
|
|
@ -537,6 +602,11 @@ u64 FileEnvironment::ReadInstruction(u32 address) {
|
|||
return code[(address - read_lowest) / sizeof(u64)];
|
||||
}
|
||||
|
||||
u32 FileEnvironment::ReadCbufSize(u32 cbuf_index) {
|
||||
const auto it = cbuf_values.find(MakeCbufKey(cbuf_index, UINT32_MAX));
|
||||
return it != cbuf_values.end() ? it->second : 0;
|
||||
}
|
||||
|
||||
u32 FileEnvironment::ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) {
|
||||
const auto it{cbuf_values.find(MakeCbufKey(cbuf_index, cbuf_offset))};
|
||||
if (it == cbuf_values.end()) {
|
||||
|
|
@ -666,11 +736,17 @@ void LoadPipelines(
|
|||
}
|
||||
u32 num_envs{};
|
||||
file.read(reinterpret_cast<char*>(&num_envs), sizeof(num_envs));
|
||||
if (num_envs == 0 || num_envs > 6) {
|
||||
throw std::ios_base::failure("Invalid shader cache stage count");
|
||||
}
|
||||
std::vector<FileEnvironment> envs(num_envs);
|
||||
for (FileEnvironment& env : envs) {
|
||||
env.Deserialize(file);
|
||||
}
|
||||
if (envs.front().ShaderStage() == Shader::Stage::Compute) {
|
||||
if (envs.size() != 1) {
|
||||
throw std::ios_base::failure("Invalid compute shader cache stage count");
|
||||
}
|
||||
load_compute(file, std::move(envs.front()));
|
||||
} else {
|
||||
load_graphics(file, std::move(envs));
|
||||
|
|
|
|||
|
|
@ -70,6 +70,9 @@ public:
|
|||
return !cbuf_replacements.empty();
|
||||
}
|
||||
|
||||
// Banks whose bound sizes were actually queried during shader translation.
|
||||
[[nodiscard]] u32 CbufSizeMask() const noexcept;
|
||||
|
||||
protected:
|
||||
std::optional<u64> TryFindSize();
|
||||
|
||||
|
|
@ -117,6 +120,8 @@ public:
|
|||
|
||||
u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override;
|
||||
|
||||
u32 ReadCbufSize(u32 cbuf_index) override;
|
||||
|
||||
Shader::TextureType ReadTextureType(u32 handle) override;
|
||||
|
||||
Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override;
|
||||
|
|
@ -143,6 +148,8 @@ public:
|
|||
|
||||
u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override;
|
||||
|
||||
u32 ReadCbufSize(u32 cbuf_index) override;
|
||||
|
||||
Shader::TextureType ReadTextureType(u32 handle) override;
|
||||
|
||||
Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override;
|
||||
|
|
@ -177,6 +184,10 @@ public:
|
|||
|
||||
[[nodiscard]] u32 ReadCbufValue(u32 cbuf_index, u32 cbuf_offset) override;
|
||||
|
||||
[[nodiscard]] u32 ReadCbufSize(u32 cbuf_index) override;
|
||||
|
||||
[[nodiscard]] u32 CbufSizeMask() const noexcept;
|
||||
|
||||
[[nodiscard]] Shader::TextureType ReadTextureType(u32 handle) override;
|
||||
|
||||
[[nodiscard]] Shader::TexturePixelFormat ReadTexturePixelFormat(u32 handle) override;
|
||||
|
|
|
|||
|
|
@ -8,6 +8,7 @@
|
|||
#include "yuzu_video_core/surface.h"
|
||||
#include "yuzu_video_core/texture_cache/format_lookup_table.h"
|
||||
#include "yuzu_video_core/texture_cache/image_info.h"
|
||||
#include "yuzu_video_core/texture_cache/mip_levels.h"
|
||||
#include "yuzu_video_core/texture_cache/samples_helper.h"
|
||||
#include "yuzu_video_core/texture_cache/types.h"
|
||||
#include "yuzu_video_core/texture_cache/util.h"
|
||||
|
|
@ -122,6 +123,9 @@ ImageInfo::ImageInfo(const TICEntry& config) noexcept {
|
|||
GetFormatType(format) != SurfaceType::ColorTexture;
|
||||
downscaleable = size.height > DownscaleHeightThreshold;
|
||||
}
|
||||
// Preserve the guest layer stride above, including any excess mip levels
|
||||
// declared by the descriptor, before limiting the host image's mip chain.
|
||||
resources.levels = MipLevelCount(config);
|
||||
}
|
||||
|
||||
ImageInfo::ImageInfo(const Maxwell3D::Regs::RenderTargetConfig& ct,
|
||||
|
|
|
|||
|
|
@ -5,6 +5,7 @@
|
|||
|
||||
#include "yuzu_common/yuzu_assert.h"
|
||||
#include "yuzu_video_core/texture_cache/image_view_info.h"
|
||||
#include "yuzu_video_core/texture_cache/mip_levels.h"
|
||||
#include "yuzu_video_core/texture_cache/texture_cache_base.h"
|
||||
#include "yuzu_video_core/texture_cache/types.h"
|
||||
#include "yuzu_video_core/texture_cache/util.h"
|
||||
|
|
@ -30,11 +31,14 @@ ImageViewInfo::ImageViewInfo(const TICEntry& config, s32 base_layer) noexcept
|
|||
: format{PixelFormatFromTIC(config)}, x_source{CastSwizzle(config.x_source)},
|
||||
y_source{CastSwizzle(config.y_source)}, z_source{CastSwizzle(config.z_source)},
|
||||
w_source{CastSwizzle(config.w_source)} {
|
||||
const s32 last_level = static_cast<s32>(MipLevelCount(config)) - 1;
|
||||
range.base = SubresourceBase{
|
||||
.level = static_cast<s32>(config.res_min_mip_level),
|
||||
.level = std::min(static_cast<s32>(config.res_min_mip_level), last_level),
|
||||
.layer = base_layer,
|
||||
};
|
||||
range.extent.levels = config.res_max_mip_level - config.res_min_mip_level + 1;
|
||||
const s32 max_level = std::clamp(static_cast<s32>(config.res_max_mip_level),
|
||||
range.base.level, last_level);
|
||||
range.extent.levels = max_level - range.base.level + 1;
|
||||
|
||||
switch (config.texture_type) {
|
||||
case TextureType::Texture1D:
|
||||
|
|
|
|||
|
|
@ -0,0 +1,33 @@
|
|||
// SPDX-License-Identifier: GPL-2.0-or-later
|
||||
|
||||
#pragma once
|
||||
|
||||
#include <algorithm>
|
||||
#include <bit>
|
||||
|
||||
#include "yuzu_video_core/textures/texture.h"
|
||||
|
||||
namespace VideoCommon {
|
||||
|
||||
// Array layers do not contribute to mip dimensions. Linear, buffer and
|
||||
// multisampled images cannot have a mip chain in the host API.
|
||||
[[nodiscard]] constexpr u32 MipLevelCount(const Tegra::Texture::TICEntry& config) {
|
||||
using Tegra::Texture::TextureType;
|
||||
if (config.IsPitchLinear() || config.IsBuffer() ||
|
||||
config.texture_type == TextureType::Texture1DBuffer ||
|
||||
config.msaa_mode != Tegra::Texture::MsaaMode::Msaa1x1) {
|
||||
return 1;
|
||||
}
|
||||
u32 max_dimension = config.Width();
|
||||
if (config.texture_type != TextureType::Texture1D &&
|
||||
config.texture_type != TextureType::Texture1DArray) {
|
||||
max_dimension = std::max(max_dimension, config.Height());
|
||||
}
|
||||
if (config.texture_type == TextureType::Texture3D) {
|
||||
max_dimension = std::max(max_dimension, config.Depth());
|
||||
}
|
||||
return std::min(config.max_mip_level.Value() + 1,
|
||||
static_cast<u32>(std::bit_width(max_dimension)));
|
||||
}
|
||||
|
||||
} // namespace VideoCommon
|
||||
|
|
@ -10,7 +10,8 @@
|
|||
namespace VideoCommon {
|
||||
|
||||
constexpr size_t NUM_RT = 8;
|
||||
constexpr size_t MAX_MIP_LEVELS = 14;
|
||||
// The TIC mip level field is four bits and encodes up to 16 levels.
|
||||
constexpr size_t MAX_MIP_LEVELS = 16;
|
||||
|
||||
constexpr Common::SlotId CORRUPT_ID{0xfffffffe};
|
||||
|
||||
|
|
|
|||
|
|
@ -31,6 +31,7 @@ VK_DEFINE_HANDLE(VmaAllocator)
|
|||
FEATURE(KHR, VariablePointer, VARIABLE_POINTERS, variable_pointer)
|
||||
|
||||
#define FOR_EACH_VK_FEATURE_1_2(FEATURE) \
|
||||
FEATURE(EXT, DescriptorIndexing, DESCRIPTOR_INDEXING, descriptor_indexing) \
|
||||
FEATURE(EXT, HostQueryReset, HOST_QUERY_RESET, host_query_reset) \
|
||||
FEATURE(KHR, 8BitStorage, 8BIT_STORAGE, bit8_storage) \
|
||||
FEATURE(KHR, TimelineSemaphore, TIMELINE_SEMAPHORE, timeline_semaphore)
|
||||
|
|
@ -337,6 +338,11 @@ public:
|
|||
return GetDriverID() != VK_DRIVER_ID_QUALCOMM_PROPRIETARY;
|
||||
}
|
||||
|
||||
/// Returns true if sampled texture array indices may vary between shader invocations.
|
||||
bool IsSampledImageArrayNonUniformIndexingSupported() const {
|
||||
return features.descriptor_indexing.shaderSampledImageArrayNonUniformIndexing != VK_FALSE;
|
||||
}
|
||||
|
||||
/// Returns true if the device supports float64 natively.
|
||||
bool IsFloat64Supported() const {
|
||||
return features.features.shaderFloat64;
|
||||
|
|
@ -377,6 +383,18 @@ public:
|
|||
return MoltenVK::PushDescriptorLimit(GetDriverID(), properties.push_descriptor.maxPushDescriptors);
|
||||
}
|
||||
|
||||
const VkPhysicalDeviceLimits& DescriptorLimits() const noexcept {
|
||||
return properties.properties.limits;
|
||||
}
|
||||
|
||||
/// Single-stage upper bound; the complete pipeline layout checks actual resource usage.
|
||||
u32 MaxBindlessDescriptorsPerStage() const {
|
||||
const auto& limits = DescriptorLimits();
|
||||
return std::min({limits.maxPerStageDescriptorSampledImages,
|
||||
limits.maxPerStageDescriptorSamplers, limits.maxDescriptorSetSampledImages,
|
||||
limits.maxDescriptorSetSamplers, limits.maxPerStageResources});
|
||||
}
|
||||
|
||||
/// Returns true if formatless image load is supported.
|
||||
bool IsFormatlessImageLoadSupported() const {
|
||||
return features.features.shaderStorageImageReadWithoutFormat;
|
||||
|
|
|
|||
|
|
@ -153,6 +153,10 @@
|
|||
<ClInclude Include="renderer_vulkan\vk_buffer_cache.h" />
|
||||
<ClInclude Include="renderer_vulkan\vk_command_pool.h" />
|
||||
<ClInclude Include="renderer_vulkan\vk_compute_pass.h" />
|
||||
<ClInclude Include="renderer_vulkan\descriptor_lookup.h" />
|
||||
<ClInclude Include="renderer_vulkan\descriptor_limits.h" />
|
||||
<ClInclude Include="renderer_vulkan\descriptor_payload.h" />
|
||||
<ClInclude Include="renderer_vulkan\cbuf_size_dependencies.h" />
|
||||
<ClInclude Include="renderer_vulkan\vk_compute_pipeline.h" />
|
||||
<ClInclude Include="renderer_vulkan\vk_descriptor_pool.h" />
|
||||
<ClInclude Include="renderer_vulkan\vk_fence_manager.h" />
|
||||
|
|
@ -192,6 +196,7 @@
|
|||
<ClInclude Include="texture_cache\format_lookup_table.h" />
|
||||
<ClInclude Include="texture_cache\image_base.h" />
|
||||
<ClInclude Include="texture_cache\image_info.h" />
|
||||
<ClInclude Include="texture_cache\mip_levels.h" />
|
||||
<ClInclude Include="texture_cache\image_view_base.h" />
|
||||
<ClInclude Include="texture_cache\image_view_info.h" />
|
||||
<ClInclude Include="texture_cache\render_targets.h" />
|
||||
|
|
|
|||
|
|
@ -474,6 +474,18 @@
|
|||
<ClInclude Include="renderer_vulkan\vk_compute_pass.h">
|
||||
<Filter>Header Files\renderer_vulkan</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="renderer_vulkan\descriptor_limits.h">
|
||||
<Filter>Header Files\renderer_vulkan</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="renderer_vulkan\descriptor_payload.h">
|
||||
<Filter>Header Files\renderer_vulkan</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="renderer_vulkan\descriptor_lookup.h">
|
||||
<Filter>Header Files\renderer_vulkan</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="renderer_vulkan\cbuf_size_dependencies.h">
|
||||
<Filter>Header Files\renderer_vulkan</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="renderer_vulkan\vk_compute_pipeline.h">
|
||||
<Filter>Header Files\renderer_vulkan</Filter>
|
||||
</ClInclude>
|
||||
|
|
@ -579,6 +591,9 @@
|
|||
<ClInclude Include="texture_cache\image_info.h">
|
||||
<Filter>Header Files\texture_cache</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="texture_cache\mip_levels.h">
|
||||
<Filter>Header Files\texture_cache</Filter>
|
||||
</ClInclude>
|
||||
<ClInclude Include="texture_cache\image_view_base.h">
|
||||
<Filter>Header Files\texture_cache</Filter>
|
||||
</ClInclude>
|
||||
|
|
|
|||
Loading…
Reference in New Issue