shifathosn
/
nxemu
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

Improve rendering compatibility on macOS (#383) 

Handle shaders the host cannot run as they are.
This commit is contained in:
Bruno 2026-10-09 01:31:06 +02:00 committed by GitHub
parent 3a25126d4f
commit 243dac50fc
No known key found for this signature in database
GPG Key ID: B5690EEEBB952194
31 changed files with 2469 additions and 49 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

4
src/yuzu_shader_recompiler/CMakeLists.txt
View File

@ -229,10 +229,13 @@ add_library(shader_recompiler STATIC
ir_opt/passes.h
ir_opt/position_pass.cpp
ir_opt/rescaling_pass.cpp
ir_opt/geometry_layer_forwarding_pass.cpp
ir_opt/shared_memory_packing_pass.cpp
ir_opt/ssa_rewrite_pass.cpp
ir_opt/texture_pass.cpp
ir_opt/vendor_workaround_pass.cpp
ir_opt/verification_pass.cpp
ir_opt/vtg_to_compute_pass.cpp
object_pool.h
precompiled_headers.h
profile.h
@ -240,6 +243,7 @@ add_library(shader_recompiler STATIC
runtime_info.h
shader_info.h
varying_state.h
vtg_as_compute.h
)
target_include_directories(shader_recompiler PRIVATE

3
src/yuzu_shader_recompiler/backend/spirv/emit_spirv_atomic.cpp
View File

@ -10,6 +10,7 @@
namespace Shader::Backend::SPIRV {
namespace {
Id SharedPointer(EmitContext& ctx, Id offset, u32 index_offset = 0) {
offset = ctx.RemapSharedMemoryOffset(offset);
const Id shift_id{ctx.Const(2U)};
Id index{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift_id)};
if (index_offset > 0) {
@ -129,12 +130,14 @@ Id EmitSharedAtomicUMax32(EmitContext& ctx, Id offset, Id value) {
}
Id EmitSharedAtomicInc32(EmitContext& ctx, Id offset, Id value) {
offset = ctx.RemapSharedMemoryOffset(offset);
const Id shift_id{ctx.Const(2U)};
const Id index{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift_id)};
return ctx.OpFunctionCall(ctx.U32[1], ctx.increment_cas_shared, index, value);
}
Id EmitSharedAtomicDec32(EmitContext& ctx, Id offset, Id value) {
offset = ctx.RemapSharedMemoryOffset(offset);
const Id shift_id{ctx.Const(2U)};
const Id index{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift_id)};
return ctx.OpFunctionCall(ctx.U32[1], ctx.decrement_cas_shared, index, value);

15
src/yuzu_shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
View File

@ -40,6 +40,21 @@ struct OutAttr {
Id type{};
};
std::optional<OutAttr> OutputAttrPointer(EmitContext& ctx, IR::Attribute attr);
} // Anonymous namespace
void EmitForwardLayer(EmitContext& ctx, IR::Attribute source) {
const u32 index{IR::GenericAttributeIndex(source)};
const u32 element{IR::GenericAttributeElement(source)};
if (!Sirit::ValidId(ctx.layer) || !Sirit::ValidId(ctx.output_generics.at(index).at(element).id)) {
return;
}
const std::optional<OutAttr> pointer{OutputAttrPointer(ctx, source)};
const Id value{ctx.OpLoad(ctx.F32[1], pointer->pointer)};
ctx.OpStore(ctx.layer, ctx.OpBitcast(ctx.U32[1], value));
}
namespace {
std::optional<OutAttr> OutputAttrPointer(EmitContext& ctx, IR::Attribute attr) {
if (IR::IsGeneric(attr)) {
const u32 index{IR::GenericAttributeIndex(attr)};

1
src/yuzu_shader_recompiler/backend/spirv/emit_spirv_instructions.h
View File

@ -34,6 +34,7 @@ void EmitWorkgroupMemoryBarrier(EmitContext& ctx);
void EmitDeviceMemoryBarrier(EmitContext& ctx);
void EmitPrologue(EmitContext& ctx);
void EmitEpilogue(EmitContext& ctx);
void EmitForwardLayer(EmitContext& ctx, IR::Attribute source);
void EmitEmitVertex(EmitContext& ctx, const IR::Value& stream);
void EmitEndPrimitive(EmitContext& ctx, const IR::Value& stream);
void EmitGetRegister(EmitContext& ctx);

29
src/yuzu_shader_recompiler/backend/spirv/emit_spirv_shared_memory.cpp
View File

@ -1,6 +1,8 @@
// SPDX-FileCopyrightText: Copyright 2021 yuzu Emulator Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include <array>
#include "yuzu_shader_recompiler/backend/spirv/emit_spirv_instructions.h"
#include "yuzu_shader_recompiler/backend/spirv/spirv_emit_context.h"
@ -28,6 +30,7 @@ std::pair<Id, Id> ExtractArgs(EmitContext& ctx, Id offset, u32 mask, u32 count)
} // Anonymous namespace
Id EmitLoadSharedU8(EmitContext& ctx, Id offset) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{
ctx.OpAccessChain(ctx.shared_u8, ctx.shared_memory_u8, ctx.u32_zero_value, offset)};
@ -39,6 +42,7 @@ Id EmitLoadSharedU8(EmitContext& ctx, Id offset) {
}
Id EmitLoadSharedS8(EmitContext& ctx, Id offset) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{
ctx.OpAccessChain(ctx.shared_u8, ctx.shared_memory_u8, ctx.u32_zero_value, offset)};
@ -50,6 +54,7 @@ Id EmitLoadSharedS8(EmitContext& ctx, Id offset) {
}
Id EmitLoadSharedU16(EmitContext& ctx, Id offset) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u16, ctx.shared_memory_u16, offset, 1)};
return ctx.OpUConvert(ctx.U32[1], ctx.OpLoad(ctx.U16, pointer));
@ -60,6 +65,7 @@ Id EmitLoadSharedU16(EmitContext& ctx, Id offset) {
}
Id EmitLoadSharedS16(EmitContext& ctx, Id offset) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u16, ctx.shared_memory_u16, offset, 1)};
return ctx.OpSConvert(ctx.U32[1], ctx.OpLoad(ctx.U16, pointer));
@ -70,6 +76,7 @@ Id EmitLoadSharedS16(EmitContext& ctx, Id offset) {
}
Id EmitLoadSharedU32(EmitContext& ctx, Id offset) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u32, ctx.shared_memory_u32, offset, 2)};
return ctx.OpLoad(ctx.U32[1], pointer);
@ -79,6 +86,7 @@ Id EmitLoadSharedU32(EmitContext& ctx, Id offset) {
}
Id EmitLoadSharedU64(EmitContext& ctx, Id offset) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u32x2, ctx.shared_memory_u32x2, offset, 3)};
return ctx.OpLoad(ctx.U32[2], pointer);
@ -94,6 +102,22 @@ Id EmitLoadSharedU64(EmitContext& ctx, Id offset) {
}
Id EmitLoadSharedU128(EmitContext& ctx, Id offset) {
if (ctx.shared_memory_16_to_12_packing) {
// The packing proof guarantees 16-byte alignment and that W is never observed. Loading
// only XYZ avoids touching the removed padding lane (which now aliases the next guest
// lane after compaction).
offset = ctx.RemapSharedMemoryOffset(offset);
const Id shift_id{ctx.Const(2U)};
const Id base_index{ctx.OpShiftRightArithmetic(ctx.U32[1], offset, shift_id)};
std::array<Id, 4> values{};
for (u32 i = 0; i < 3; ++i) {
const Id index{i == 0 ? base_index : ctx.OpIAdd(ctx.U32[1], base_index, ctx.Const(i))};
const Id pointer{ctx.OpAccessChain(ctx.shared_u32, ctx.shared_memory_u32, index)};
values[i] = ctx.OpLoad(ctx.U32[1], pointer);
}
values[3] = ctx.u32_zero_value;
return ctx.OpCompositeConstruct(ctx.U32[4], values);
}
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u32x4, ctx.shared_memory_u32x4, offset, 4)};
return ctx.OpLoad(ctx.U32[4], pointer);
@ -110,6 +134,7 @@ Id EmitLoadSharedU128(EmitContext& ctx, Id offset) {
}
void EmitWriteSharedU8(EmitContext& ctx, Id offset, Id value) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{
ctx.OpAccessChain(ctx.shared_u8, ctx.shared_memory_u8, ctx.u32_zero_value, offset)};
@ -120,6 +145,7 @@ void EmitWriteSharedU8(EmitContext& ctx, Id offset, Id value) {
}
void EmitWriteSharedU16(EmitContext& ctx, Id offset, Id value) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u16, ctx.shared_memory_u16, offset, 1)};
ctx.OpStore(pointer, ctx.OpUConvert(ctx.U16, value));
@ -129,6 +155,7 @@ void EmitWriteSharedU16(EmitContext& ctx, Id offset, Id value) {
}
void EmitWriteSharedU32(EmitContext& ctx, Id offset, Id value) {
offset = ctx.RemapSharedMemoryOffset(offset);
Id pointer{};
if (ctx.profile.support_explicit_workgroup_layout) {
pointer = Pointer(ctx, ctx.shared_u32, ctx.shared_memory_u32, offset, 2);
@ -141,6 +168,7 @@ void EmitWriteSharedU32(EmitContext& ctx, Id offset, Id value) {
}
void EmitWriteSharedU64(EmitContext& ctx, Id offset, Id value) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u32x2, ctx.shared_memory_u32x2, offset, 3)};
ctx.OpStore(pointer, value);
@ -156,6 +184,7 @@ void EmitWriteSharedU64(EmitContext& ctx, Id offset, Id value) {
}
void EmitWriteSharedU128(EmitContext& ctx, Id offset, Id value) {
offset = ctx.RemapSharedMemoryOffset(offset);
if (ctx.profile.support_explicit_workgroup_layout) {
const Id pointer{Pointer(ctx, ctx.shared_u32x4, ctx.shared_memory_u32x4, offset, 4)};
ctx.OpStore(pointer, value);

3
src/yuzu_shader_recompiler/backend/spirv/emit_spirv_special.cpp
View File

@ -125,6 +125,9 @@ void EmitEpilogue(EmitContext& ctx) {
!ctx.profile.support_native_ndc) {
ConvertDepthMode(ctx);
}
if (ctx.stage == Stage::VertexB && ctx.runtime_info.forwarded_layer_attribute) {
EmitForwardLayer(ctx, *ctx.runtime_info.forwarded_layer_attribute);
}
if (ctx.stage == Stage::Fragment) {
AlphaTest(ctx);
}

24
src/yuzu_shader_recompiler/backend/spirv/spirv_emit_context.cpp
View File

@ -461,6 +461,7 @@ EmitContext::EmitContext(const Profile& profile_, const RuntimeInfo& runtime_inf
: Sirit::Module(profile_.supported_spirv), profile{profile_}, runtime_info{runtime_info_},
stage{program.stage}, texture_rescaling_index{bindings.texture_scaling_index},
image_rescaling_index{bindings.image_scaling_index} {
shared_memory_16_to_12_packing = program.shared_memory_16_to_12_packing;
const bool is_unified{profile.unified_descriptor_binding};
u32& uniform_binding{is_unified ? bindings.unified : bindings.uniform_buffer};
u32& storage_binding{is_unified ? bindings.unified : bindings.storage_buffer};
@ -527,6 +528,18 @@ Id EmitContext::BitOffset16(const IR::Value& offset) {
return OpBitwiseAnd(U32[1], OpShiftLeftLogical(U32[1], Def(offset), Const(3u)), Const(16u));
}
Id EmitContext::RemapSharedMemoryOffset(Id offset) {
if (!shared_memory_16_to_12_packing) {
return offset;
}
// Compact each 16-byte guest lane to the 12 bytes proven live by
// TryPackSharedMemory16To12(). offset - (offset / 16) * 4 is injective for
// offsets 0..11 within every source lane.
const Id lane{OpShiftRightLogical(U32[1], offset, Const(4U))};
const Id removed_padding{OpShiftLeftLogical(U32[1], lane, Const(2U))};
return OpISub(U32[1], offset, removed_padding);
}
void EmitContext::DefineCommonTypes(const Info& info) {
void_id = TypeVoid();
@ -598,8 +611,15 @@ void EmitContext::DefineSharedMemory(const IR::Program& program) {
if (program.shared_memory_size == 0) {
return;
}
u32 shared_memory_size{program.shared_memory_size};
const u32 max_size{profile.max_compute_shared_memory_size};
if (max_size != 0 && shared_memory_size > max_size) {
LOG_WARNING(Shader_SPIRV, "Requested shared memory size ({}) exceeds device limit ({})",
shared_memory_size, max_size);
shared_memory_size = max_size;
}
const auto make{[&](Id element_type, u32 element_size) {
const u32 num_elements{Common::DivCeil(program.shared_memory_size, element_size)};
const u32 num_elements{Common::DivCeil(shared_memory_size, element_size)};
const Id array_type{TypeArray(element_type, Const(num_elements))};
Decorate(array_type, spv::Decoration::ArrayStride, element_size);
@ -634,7 +654,7 @@ void EmitContext::DefineSharedMemory(const IR::Program& program) {
std::tie(shared_memory_u32x4, shared_u32x4, std::ignore) = make(U32[4], 16);
return;
}
const u32 num_elements{Common::DivCeil(program.shared_memory_size, 4U)};
const u32 num_elements{Common::DivCeil(shared_memory_size, 4U)};
const Id type{TypeArray(U32[1], Const(num_elements))};
shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type);

2
src/yuzu_shader_recompiler/backend/spirv/spirv_emit_context.h
View File

@ -163,6 +163,7 @@ public:
[[nodiscard]] Id BitOffset8(const IR::Value& offset);
[[nodiscard]] Id BitOffset16(const IR::Value& offset);
[[nodiscard]] Id RemapSharedMemoryOffset(Id offset);
Id Const(u32 value) {
return Constant(U32[1], value);
@ -302,6 +303,7 @@ public:
Id local_memory{};
bool shared_memory_16_to_12_packing{};
Id shared_memory_u8{};
Id shared_memory_u16{};
Id shared_memory_u32{};

1
src/yuzu_shader_recompiler/frontend/ir/program.h
View File

@ -26,6 +26,7 @@ struct Program {
u32 invocations{};
u32 local_memory_size{};
u32 shared_memory_size{};
bool shared_memory_16_to_12_packing{};
bool is_geometry_passthrough{};
};

71
src/yuzu_shader_recompiler/frontend/maxwell/translate_program.cpp
View File

@ -435,4 +435,75 @@ IR::Program GenerateGeometryPassthrough(ObjectPool<IR::Inst>& inst_pool,
return program;
}
IR::Program GenerateVtgPassthroughVertex(ObjectPool<IR::Inst>& inst_pool,
ObjectPool<IR::Block>& block_pool,
const IR::Program& source_program,
const VaryingState& source_stores,
const VtgVaryingLayout& layout) {
static constexpr u32 word_size{sizeof(u32)};
IR::Program program;
program.stage = Stage::VertexB;
program.info.legacy_stores_mapping = source_program.info.legacy_stores_mapping;
program.info.used_clip_distances = source_program.info.used_clip_distances;
program.info.loads.Set(IR::Attribute::VertexId);
program.info.used_storage_buffer_types = IR::Type::U32;
program.info.storage_buffers_descriptors.push_back({
.cbuf_index = static_cast<u32>(Info::MAX_CBUFS),
.cbuf_offset = 0,
.count = 1,
.is_written = false,
});
IR::Block* const current_block{block_pool.Create(inst_pool)};
auto& node{program.syntax_list.emplace_back()};
node.type = IR::AbstractSyntaxNode::Type::Block;
node.data.block = current_block;
IR::IREmitter ir{*current_block};
ir.Prologue();
const IR::U32 vertex{ir.GetAttributeU32(IR::Attribute::VertexId)};
const IR::U32 base{ir.IMul(vertex, ir.Imm32(layout.stride * word_size))};
const auto forward{[&](IR::Attribute attribute) {
const std::optional<u32> word{layout.Offset(attribute)};
if (!word) {
return;
}
const IR::U32 offset{ir.IAdd(base, ir.Imm32(*word * word_size))};
current_block->AppendNewInst(IR::Opcode::LoadStorage32,
{IR::Value{0U}, IR::Value{offset}});
const IR::U32 value{IR::Value{&current_block->back()}};
ir.SetAttribute(attribute, ir.BitCast<IR::F32>(value), ir.Imm32(0U));
program.info.stores.Set(attribute);
}};
for (u32 element = 0; element < 4; ++element) {
forward(IR::Attribute::PositionX + element);
}
for (u32 index = 0; index < IR::NUM_GENERICS; ++index) {
for (u32 element = 0; element < 4; ++element) {
forward(IR::Attribute::Generic0X + index * 4 + element);
}
}
for (u32 index = 0; index < 8; ++index) {
if (source_stores[IR::Attribute::ClipDistance0 + index]) {
forward(IR::Attribute::ClipDistance0 + index);
}
}
forward(IR::Attribute::PointSize);
forward(IR::Attribute::Layer);
forward(IR::Attribute::ViewportIndex);
IR::Block* const return_block{block_pool.Create(inst_pool)};
IR::IREmitter{*return_block}.Epilogue();
current_block->AddBranch(return_block);
auto& merge{program.syntax_list.emplace_back()};
merge.type = IR::AbstractSyntaxNode::Type::Block;
merge.data.block = return_block;
program.syntax_list.emplace_back().type = IR::AbstractSyntaxNode::Type::Return;
program.blocks = GenerateBlocks(program.syntax_list);
program.post_order_blocks = PostOrder(program.syntax_list.front());
return program;
}
} // namespace Shader::Maxwell

9
src/yuzu_shader_recompiler/frontend/maxwell/translate_program.h
View File

@ -6,6 +6,7 @@
#include "yuzu_shader_recompiler/environment.h"
#include "yuzu_shader_recompiler/frontend/ir/basic_block.h"
#include "yuzu_shader_recompiler/frontend/ir/program.h"
#include "yuzu_shader_recompiler/vtg_as_compute.h"
#include "yuzu_shader_recompiler/frontend/maxwell/control_flow.h"
#include "yuzu_shader_recompiler/object_pool.h"
#include "yuzu_shader_recompiler/runtime_info.h"
@ -34,4 +35,12 @@ void ConvertLegacyToGeneric(IR::Program& program, const RuntimeInfo& runtime_inf
IR::Program& source_program,
Shader::OutputTopology output_topology);
// Creates the vertex shader feeding the rasterizer with the vertices written to a storage buffer
// by the stages running as compute, indexed by the vertex index of the draw.
[[nodiscard]] IR::Program GenerateVtgPassthroughVertex(ObjectPool<IR::Inst>& inst_pool,
ObjectPool<IR::Block>& block_pool,
const IR::Program& source_program,
const VaryingState& source_stores,
const VtgVaryingLayout& layout);
} // namespace Shader::Maxwell

112
src/yuzu_shader_recompiler/ir_opt/geometry_layer_forwarding_pass.cpp Normal file
View File

@ -0,0 +1,112 @@
// SPDX-FileCopyrightText: Copyright 2026 NxEmu Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include <array>
#include <optional>
#include "yuzu_shader_recompiler/frontend/ir/attribute.h"
#include "yuzu_shader_recompiler/frontend/ir/basic_block.h"
#include "yuzu_shader_recompiler/frontend/ir/program.h"
#include "yuzu_shader_recompiler/frontend/ir/value.h"
#include "yuzu_shader_recompiler/ir_opt/passes.h"
namespace Shader::Optimization {
namespace {
constexpr size_t MAX_VERTICES = 3;
bool HasSideEffects(IR::Opcode opcode) {
switch (opcode) {
case IR::Opcode::SetAttributeIndexed:
case IR::Opcode::SetPatch:
case IR::Opcode::WriteGlobalU8:
case IR::Opcode::WriteGlobalS8:
case IR::Opcode::WriteGlobalU16:
case IR::Opcode::WriteGlobalS16:
case IR::Opcode::WriteGlobal32:
case IR::Opcode::WriteGlobal64:
case IR::Opcode::WriteGlobal128:
case IR::Opcode::WriteStorageU8:
case IR::Opcode::WriteStorageS8:
case IR::Opcode::WriteStorageU16:
case IR::Opcode::WriteStorageS16:
case IR::Opcode::WriteStorage32:
case IR::Opcode::WriteStorage64:
case IR::Opcode::WriteStorage128:
case IR::Opcode::ImageWrite:
case IR::Opcode::BindlessImageWrite:
case IR::Opcode::BoundImageWrite:
return true;
default:
return false;
}
}
} // Anonymous namespace
std::optional<IR::Attribute> FindForwardedLayerAttribute(const IR::Program& program,
u32 input_vertices) {
if (program.stage != Stage::Geometry || program.is_geometry_passthrough ||
program.invocations > 1 || input_vertices != MAX_VERTICES ||
program.output_vertices != input_vertices ||
program.output_topology != OutputTopology::TriangleStrip) {
return std::nullopt;
}
std::optional<IR::Attribute> layer_source;
std::array<IR::Value, MAX_VERTICES> emitted_vertices;
std::optional<IR::Value> current_vertex;
size_t num_emitted{};
for (const IR::Block* const block : program.blocks) {
for (const IR::Inst& inst : block->Instructions()) {
const IR::Opcode opcode{inst.GetOpcode()};
if (HasSideEffects(opcode)) {
return std::nullopt;
}
if (opcode == IR::Opcode::EmitVertex) {
if (!current_vertex || num_emitted == MAX_VERTICES) {
return std::nullopt;
}
for (size_t index = 0; index < num_emitted; ++index) {
if (emitted_vertices[index] == *current_vertex) {
return std::nullopt;
}
}
emitted_vertices[num_emitted++] = *current_vertex;
current_vertex.reset();
continue;
}
if (opcode != IR::Opcode::SetAttribute) {
continue;
}
// Every output has to be the same attribute of one input vertex
const IR::Attribute target{inst.Arg(0).Attribute()};
const IR::Value value{inst.Arg(1).Resolve()};
if (value.IsImmediate()) {
return std::nullopt;
}
const IR::Inst* const source{value.InstRecursive()};
if (source->GetOpcode() != IR::Opcode::GetAttribute) {
return std::nullopt;
}
const IR::Attribute source_attribute{source->Arg(0).Attribute()};
const IR::Value vertex{source->Arg(1).Resolve()};
if (current_vertex && !(*current_vertex == vertex)) {
return std::nullopt;
}
current_vertex = vertex;
if (target == IR::Attribute::Layer) {
if (!IR::IsGeneric(source_attribute) ||
(layer_source && *layer_source != source_attribute)) {
return std::nullopt;
}
layer_source = source_attribute;
} else if (target != source_attribute) {
return std::nullopt;
}
}
}
if (num_emitted != input_vertices || current_vertex) {
return std::nullopt;
}
return layer_source;
}
} // namespace Shader::Optimization

22
src/yuzu_shader_recompiler/ir_opt/passes.h
View File

@ -3,8 +3,12 @@
#pragma once
#include <optional>
#include "yuzu_shader_recompiler/frontend/ir/attribute.h"
#include "yuzu_shader_recompiler/environment.h"
#include "yuzu_shader_recompiler/frontend/ir/program.h"
#include "yuzu_shader_recompiler/vtg_as_compute.h"
namespace Shader {
struct HostTranslateInfo;
@ -25,6 +29,24 @@ void RescalingPass(IR::Program& program);
void SsaRewritePass(IR::Program& program);
void PositionPass(Environment& env, IR::Program& program);
void TexturePass(Environment& env, IR::Program& program, const HostTranslateInfo& host_info);
// SPIR-V only: compact oversized compute shared memory when every live access fits XYZ.
// Preserves valid guest accesses; does not prove or repair guest out-of-bounds addresses.
// On failure the program is unchanged. Run after optimization and before SPIR-V emission.
[[nodiscard]] bool TryPackSharedMemory16To12(IR::Program& program, u32 max_shared_memory);
// Returns the input attribute a geometry shader copies into the layer when all it does is
// forward its input triangle unchanged and select the layer, so the previous stage can do it.
[[nodiscard]] std::optional<IR::Attribute> FindForwardedLayerAttribute(const IR::Program& program,
u32 input_vertices);
// Turns a vertex shader into a compute shader pulling its attributes from the vertex buffers and
// writing its outputs to a storage buffer. On failure the program is unchanged.
[[nodiscard]] bool VtgVertexToCompute(IR::Program& program, const VtgVertexInputs& inputs,
const VtgVaryingLayout& layout, VtgVertexBindings& bindings);
// Turns a geometry shader into a compute shader reading the outputs of the vertex stage and
// writing its vertices and a strip index buffer. On failure the program is unchanged.
[[nodiscard]] bool VtgGeometryToCompute(IR::Program& program, VtgTopology topology,
u32 input_vertices, const VtgVaryingLayout& input_layout,
const VtgVaryingLayout& output_layout,
VtgGeometryBindings& bindings);
void LayerPass(IR::Program& program, const HostTranslateInfo& host_info);
void VendorWorkaroundPass(IR::Program& program);
void VerificationPass(const IR::Program& program);

343
src/yuzu_shader_recompiler/ir_opt/shared_memory_packing_pass.cpp Normal file
View File

@ -0,0 +1,343 @@
// SPDX-FileCopyrightText: Copyright 2026 Eden Emulator Project
// SPDX-License-Identifier: GPL-3.0-or-later
#include <algorithm>
#include <cstdint>
#include <unordered_map>
#include <unordered_set>
#include <vector>
#include "yuzu_shader_recompiler/frontend/ir/basic_block.h"
#include "yuzu_shader_recompiler/frontend/ir/opcodes.h"
#include "yuzu_shader_recompiler/frontend/ir/program.h"
#include "yuzu_shader_recompiler/frontend/ir/value.h"
#include "yuzu_shader_recompiler/ir_opt/passes.h"
namespace Shader::Optimization {
namespace {
using ResidueMask = std::uint16_t;
constexpr ResidueMask ALL_RESIDUES = 0xffff;
constexpr u32 SOURCE_STRIDE = 16;
constexpr u32 PACKED_STRIDE = 12;
[[nodiscard]] constexpr ResidueMask Residue(u32 value) {
return static_cast<ResidueMask>(1U << (value & (SOURCE_STRIDE - 1)));
}
[[nodiscard]] ResidueMask BinaryResidues(ResidueMask lhs, ResidueMask rhs,
u32 (*operation)(u32, u32)) {
ResidueMask result{};
for (u32 x = 0; x < SOURCE_STRIDE; ++x) {
if ((lhs & Residue(x)) == 0) {
continue;
}
for (u32 y = 0; y < SOURCE_STRIDE; ++y) {
if ((rhs & Residue(y)) == 0) {
continue;
}
result |= Residue(operation(x, y));
}
}
return result;
}
[[nodiscard]] u32 Add(u32 lhs, u32 rhs) {
return lhs + rhs;
}
[[nodiscard]] u32 Sub(u32 lhs, u32 rhs) {
return lhs - rhs;
}
[[nodiscard]] u32 Mul(u32 lhs, u32 rhs) {
return lhs * rhs;
}
[[nodiscard]] u32 And(u32 lhs, u32 rhs) {
return lhs & rhs;
}
[[nodiscard]] u32 Or(u32 lhs, u32 rhs) {
return lhs | rhs;
}
[[nodiscard]] u32 Xor(u32 lhs, u32 rhs) {
return lhs ^ rhs;
}
class ResidueAnalyzer {
public:
[[nodiscard]] ResidueMask Get(const IR::Value& value) {
if (!value.IsIdentity() && value.IsImmediate()) {
return value.Type() == IR::Type::U32 ? Residue(value.U32()) : ALL_RESIDUES;
}
// Bound recursion, including Identity chains, for large guest shaders.
if (visiting.size() >= 128) {
return ALL_RESIDUES;
}
const IR::Inst* const inst{value.Inst()};
if (const auto it = cache.find(inst); it != cache.end()) {
return it->second;
}
if (!visiting.insert(inst).second) {
// Be deliberately conservative for loop-carried Phi values.
return ALL_RESIDUES;
}
const ResidueMask result{Analyze(*inst)};
visiting.erase(inst);
cache.emplace(inst, result);
return result;
}
private:
[[nodiscard]] ResidueMask Analyze(const IR::Inst& inst) {
const auto binary = [&](u32 (*operation)(u32, u32)) {
return BinaryResidues(Get(inst.Arg(0)), Get(inst.Arg(1)), operation);
};
switch (inst.GetOpcode()) {
case IR::Opcode::Identity:
return Get(inst.Arg(0));
case IR::Opcode::Phi: {
if (inst.NumArgs() == 0) {
return ALL_RESIDUES;
}
ResidueMask result{};
for (size_t i = 0; i < inst.NumArgs(); ++i) {
result |= Get(inst.Arg(i));
}
return result;
}
case IR::Opcode::SelectU32:
return Get(inst.Arg(1)) | Get(inst.Arg(2));
case IR::Opcode::IAdd32:
return binary(Add);
case IR::Opcode::ISub32:
return binary(Sub);
case IR::Opcode::IMul32:
return binary(Mul);
case IR::Opcode::BitwiseAnd32:
return binary(And);
case IR::Opcode::BitwiseOr32:
return binary(Or);
case IR::Opcode::BitwiseXor32:
return binary(Xor);
case IR::Opcode::ShiftLeftLogical32: {
const IR::Value shift{inst.Arg(1)};
if (!shift.IsImmediate()) {
return ALL_RESIDUES;
}
const u32 amount{shift.U32()};
if (amount >= 32) {
return ALL_RESIDUES;
}
if (amount >= 4) {
return Residue(0);
}
ResidueMask result{};
const ResidueMask source{Get(inst.Arg(0))};
for (u32 value = 0; value < SOURCE_STRIDE; ++value) {
if ((source & Residue(value)) != 0) {
result |= Residue(value << amount);
}
}
return result;
}
default:
// Unknown sources are fine when a parent operation subsequently proves alignment
// (for example, local invocation ID multiplied by 160). Returning all low bits keeps
// this analysis safe and intentionally conservative.
return ALL_RESIDUES;
}
}
std::unordered_map<const IR::Inst*, ResidueMask> cache;
std::unordered_set<const IR::Inst*> visiting;
};
// Build direct use lists once instead of rescanning the shader for each vector load.
class SharedLoadUses {
public:
explicit SharedLoadUses(const IR::Program& program) {
for (const IR::Block* const block : program.blocks) {
for (const IR::Inst& user : *block) {
for (size_t index = 0; index < user.NumArgs(); ++index) {
const IR::Value arg{user.Arg(index)};
if (arg.IsIdentity() || !arg.IsImmediate()) {
uses[arg.Inst()].push_back({&user, index});
}
}
}
}
}
[[nodiscard]] bool OnlyUsesXYZ(const IR::Inst* load) const {
std::vector<const IR::Inst*> pending{load};
std::unordered_set<const IR::Inst*> visited;
while (!pending.empty()) {
const IR::Inst* const value{pending.back()};
pending.pop_back();
if (!visited.insert(value).second) {
return false;
}
const auto it = uses.find(value);
const size_t observed_uses{it == uses.end() ? 0 : it->second.size()};
// Also reject uses outside program.blocks (including associated/pseudo uses).
if (observed_uses != static_cast<size_t>(value->UseCount())) {
return false;
}
if (it == uses.end()) {
continue;
}
for (const auto& [user, index] : it->second) {
if (user->GetOpcode() == IR::Opcode::Identity && index == 0) {
pending.push_back(user);
continue;
}
if (user->GetOpcode() != IR::Opcode::CompositeExtractU32x4 || index != 0) {
return false;
}
const IR::Value component{user->Arg(1)};
if (!component.IsImmediate() || component.U32() >= 3) {
return false;
}
}
}
return true;
}
private:
struct Use {
const IR::Inst* user;
size_t index;
};
std::unordered_map<const IR::Inst*, std::vector<Use>> uses;
};
[[nodiscard]] bool AccessFitsPackedLane(ResidueMask residues, u32 access_size) {
if (residues == 0) {
return false;
}
for (u32 residue = 0; residue < SOURCE_STRIDE; ++residue) {
if ((residues & Residue(residue)) == 0) {
continue;
}
// The explicit backend rounds vector addresses down to their natural alignment.
// Require that alignment before switching to the scalar packed representation.
if (residue % access_size != 0 || residue + access_size > PACKED_STRIDE) {
return false;
}
}
return true;
}
[[nodiscard]] bool IsSafeAccess(const IR::Inst& inst, ResidueAnalyzer& residues,
const SharedLoadUses& uses) {
const ResidueMask offset_residues{residues.Get(inst.Arg(0))};
switch (inst.GetOpcode()) {
case IR::Opcode::LoadSharedU8:
case IR::Opcode::LoadSharedS8:
case IR::Opcode::WriteSharedU8:
return AccessFitsPackedLane(offset_residues, 1);
case IR::Opcode::LoadSharedU16:
case IR::Opcode::LoadSharedS16:
case IR::Opcode::WriteSharedU16:
return AccessFitsPackedLane(offset_residues, 2);
case IR::Opcode::LoadSharedU32:
case IR::Opcode::WriteSharedU32:
case IR::Opcode::SharedAtomicIAdd32:
case IR::Opcode::SharedAtomicSMin32:
case IR::Opcode::SharedAtomicUMin32:
case IR::Opcode::SharedAtomicSMax32:
case IR::Opcode::SharedAtomicUMax32:
case IR::Opcode::SharedAtomicInc32:
case IR::Opcode::SharedAtomicDec32:
case IR::Opcode::SharedAtomicAnd32:
case IR::Opcode::SharedAtomicOr32:
case IR::Opcode::SharedAtomicXor32:
case IR::Opcode::SharedAtomicExchange32:
return AccessFitsPackedLane(offset_residues, 4);
case IR::Opcode::LoadSharedU64:
case IR::Opcode::WriteSharedU64:
case IR::Opcode::SharedAtomicExchange32x2:
return AccessFitsPackedLane(offset_residues, 8);
case IR::Opcode::SharedAtomicExchange64:
// Packing intentionally disables explicit workgroup layout. On profiles with native
// shared int64 atomics that would turn this operation into the backend's non-atomic
// 2x32 fallback, changing guest semantics. Fail closed instead.
return false;
case IR::Opcode::LoadSharedU128:
// A 16-byte guest load spans XYZ + a four-byte padding lane in the pattern this pass
// handles. We may drop W only when the IR proves it is never observed.
return offset_residues == Residue(0) && uses.OnlyUsesXYZ(&inst);
case IR::Opcode::WriteSharedU128:
// Dropping a written W component would change observable shared-memory state.
return false;
default:
return false;
}
}
[[nodiscard]] bool IsSharedAccess(IR::Opcode opcode) {
switch (opcode) {
case IR::Opcode::LoadSharedU8:
case IR::Opcode::LoadSharedS8:
case IR::Opcode::LoadSharedU16:
case IR::Opcode::LoadSharedS16:
case IR::Opcode::LoadSharedU32:
case IR::Opcode::LoadSharedU64:
case IR::Opcode::LoadSharedU128:
case IR::Opcode::WriteSharedU8:
case IR::Opcode::WriteSharedU16:
case IR::Opcode::WriteSharedU32:
case IR::Opcode::WriteSharedU64:
case IR::Opcode::WriteSharedU128:
case IR::Opcode::SharedAtomicIAdd32:
case IR::Opcode::SharedAtomicSMin32:
case IR::Opcode::SharedAtomicUMin32:
case IR::Opcode::SharedAtomicSMax32:
case IR::Opcode::SharedAtomicUMax32:
case IR::Opcode::SharedAtomicInc32:
case IR::Opcode::SharedAtomicDec32:
case IR::Opcode::SharedAtomicAnd32:
case IR::Opcode::SharedAtomicOr32:
case IR::Opcode::SharedAtomicXor32:
case IR::Opcode::SharedAtomicExchange32:
case IR::Opcode::SharedAtomicExchange64:
case IR::Opcode::SharedAtomicExchange32x2:
return true;
default:
return false;
}
}
[[nodiscard]] u32 PackedSize(u32 source_size) {
const u32 full_lanes{source_size / SOURCE_STRIDE};
const u32 tail{source_size % SOURCE_STRIDE};
const u32 bytes{full_lanes * PACKED_STRIDE + std::min(tail, PACKED_STRIDE)};
// The SPIR-V backend allocates whole u32 words, including any partial tail.
return (bytes + 3U) & ~3U;
}
} // Anonymous namespace
bool TryPackSharedMemory16To12(IR::Program& program, u32 max_shared_memory) {
if (program.stage != Stage::Compute || program.shared_memory_size <= max_shared_memory ||
program.shared_memory_16_to_12_packing) {
return false;
}
const u32 packed_size{PackedSize(program.shared_memory_size)};
if (packed_size >= program.shared_memory_size || packed_size > max_shared_memory) {
return false;
}
ResidueAnalyzer residues;
const SharedLoadUses uses{program};
for (const IR::Block* const block : program.blocks) {
for (const IR::Inst& inst : *block) {
if (IsSharedAccess(inst.GetOpcode()) && !IsSafeAccess(inst, residues, uses)) {
return false;
}
}
}
program.shared_memory_size = packed_size;
program.shared_memory_16_to_12_packing = true;
return true;
}
} // namespace Shader::Optimization

566
src/yuzu_shader_recompiler/ir_opt/vtg_to_compute_pass.cpp Normal file
View File

@ -0,0 +1,566 @@
// SPDX-FileCopyrightText: Copyright 2026 NxEmu Project
// SPDX-License-Identifier: GPL-2.0-or-later
#include <array>
#include <optional>
#include "yuzu_common/div_ceil.h"
#include "yuzu_shader_recompiler/frontend/ir/attribute.h"
#include "yuzu_shader_recompiler/frontend/ir/basic_block.h"
#include "yuzu_shader_recompiler/frontend/ir/ir_emitter.h"
#include "yuzu_shader_recompiler/frontend/ir/program.h"
#include "yuzu_shader_recompiler/frontend/ir/value.h"
#include "yuzu_shader_recompiler/ir_opt/passes.h"
#include "yuzu_shader_recompiler/vtg_as_compute.h"
namespace Shader::Optimization {
namespace {
constexpr u32 WORD_SIZE = sizeof(u32);
constexpr size_t MAX_INPUT_VERTICES = 6;
class Emitter : public IR::IREmitter {
public:
explicit Emitter(IR::Block& block_, IR::Block::iterator point_)
: IR::IREmitter{block_, point_}, point{point_} {}
IR::U32 Add(const IR::U32& a, const IR::U32& b) {
return IR::U32{IAdd(a, b)};
}
IR::U32 Sub(const IR::U32& a, const IR::U32& b) {
return IR::U32{ISub(a, b)};
}
IR::U32 Mul(const IR::U32& a, u32 b) {
return b == 1 ? a : IMul(a, Imm32(b));
}
IR::U32 Mod(const IR::U32& a, const IR::U32& b) {
return Sub(a, IMul(IDiv(a, b), b));
}
IR::U32 LoadStorage(u32 binding, const IR::U32& byte_offset) {
const auto it{block->PrependNewInst(point, IR::Opcode::LoadStorage32,
{IR::Value{binding}, IR::Value{byte_offset}})};
return IR::U32{IR::Value{&*it}};
}
void WriteStorage(u32 binding, const IR::U32& byte_offset, const IR::U32& value) {
block->PrependNewInst(point, IR::Opcode::WriteStorage32,
{IR::Value{binding}, IR::Value{byte_offset}, IR::Value{value}});
}
IR::U32 LoadWord(u32 binding, u32 word) {
return LoadStorage(binding, Imm32(word * WORD_SIZE));
}
/// Reads up to four bytes from any byte offset of a storage buffer
IR::U32 LoadUnaligned(u32 binding, const IR::U32& byte_offset, const IR::U32& num_bytes) {
const IR::U32 misalign{BitwiseAnd(byte_offset, Imm32(3U))};
const IR::U32 base{Sub(byte_offset, misalign)};
const IR::U32 shift{ShiftLeftLogical(misalign, Imm32(3U))};
const IR::U32 low{ShiftRightLogical(LoadStorage(binding, base), shift)};
// The next word is only touched by reads crossing into it, as it may not exist
const IR::U1 crosses{ILessThan(Imm32(WORD_SIZE), Add(misalign, num_bytes), false)};
const IR::U32 next{Add(base, IR::U32{Select(crosses, Imm32(WORD_SIZE), Imm32(0U))})};
const IR::U32 high_shift{BitwiseAnd(Sub(Imm32(32U), shift), Imm32(31U))};
const IR::U32 high{ShiftLeftLogical(LoadStorage(binding, next), high_shift)};
return BitwiseOr(low, IR::U32{Select(crosses, high, Imm32(0U))});
}
private:
IR::Block::iterator point;
};
IR::Block::iterator IteratorTo(IR::Inst& inst) {
return IR::Block::InstructionList::s_iterator_to(inst);
}
u32 ReserveStorageBuffer(Info& info, bool is_written) {
const u32 index{static_cast<u32>(info.storage_buffers_descriptors.size())};
info.storage_buffers_descriptors.push_back({
.cbuf_index = static_cast<u32>(Info::MAX_CBUFS),
.cbuf_offset = index,
.count = 1,
.is_written = is_written,
});
return index;
}
bool IsFloatLike(VtgComponentType type) {
return type != VtgComponentType::UInt && type != VtgComponentType::SInt;
}
/// Converts a component as read from a vertex buffer into the value the shader receives
IR::F32 DecodeComponent(Emitter& ir, const VtgVertexAttribute& attribute, const IR::U32& raw,
u32 bits) {
const auto unsigned_value{[&] {
return bits >= 32 ? raw : ir.BitFieldExtract(raw, ir.Imm32(0U), ir.Imm32(bits), false);
}};
const auto signed_value{[&] {
return bits >= 32 ? raw : ir.BitFieldExtract(raw, ir.Imm32(0U), ir.Imm32(bits), true);
}};
switch (attribute.type) {
case VtgComponentType::Float:
switch (bits) {
case 32:
return ir.BitCast<IR::F32>(raw);
case 16:
return IR::F32{ir.CompositeExtract(ir.UnpackHalf2x16(unsigned_value()), 0)};
default: {
// Small unsigned floats are half floats missing the sign and low mantissa bits
const IR::U32 half{ir.ShiftLeftLogical(unsigned_value(), ir.Imm32(15U - bits))};
return IR::F32{ir.CompositeExtract(ir.UnpackHalf2x16(half), 0)};
}
}
case VtgComponentType::UNorm: {
const f64 max_value{static_cast<f64>((u64{1} << bits) - 1)};
return IR::F32{ir.FPMul(IR::F32{ir.ConvertUToF(32, 32, unsigned_value())},
ir.Imm32(static_cast<f32>(1.0 / max_value)))};
}
case VtgComponentType::SNorm: {
const f64 max_value{bits > 1 ? static_cast<f64>((u64{1} << (bits - 1)) - 1) : 1.0};
const IR::F32 value{ir.FPMul(IR::F32{ir.ConvertSToF(32, 32, signed_value())},
ir.Imm32(static_cast<f32>(1.0 / max_value)))};
return IR::F32{ir.FPMax(value, ir.Imm32(-1.0f))};
}
case VtgComponentType::UScaled:
return IR::F32{ir.ConvertUToF(32, 32, unsigned_value())};
case VtgComponentType::SScaled:
return IR::F32{ir.ConvertSToF(32, 32, signed_value())};
case VtgComponentType::UInt:
return ir.BitCast<IR::F32>(unsigned_value());
case VtgComponentType::SInt:
return ir.BitCast<IR::F32>(signed_value());
case VtgComponentType::Disabled:
break;
}
return ir.Imm32(0.0f);
}
IR::F32 DefaultComponent(Emitter& ir, const VtgVertexAttribute& attribute, u32 component) {
if (component != 3) {
return ir.Imm32(0.0f);
}
return IsFloatLike(attribute.type) ? ir.Imm32(1.0f) : ir.BitCast<IR::F32>(ir.Imm32(1U));
}
/// Value of a varying missing from the buffer it is read from
IR::F32 MissingVarying(Emitter& ir, IR::Attribute attribute) {
const bool is_w{IR::IsGeneric(attribute) && IR::GenericAttributeElement(attribute) == 3};
return ir.Imm32(is_w ? 1.0f : 0.0f);
}
bool IsIndexedOrUnsupported(const IR::Program& program) {
const Info& info{program.info};
if (info.loads_indexed_attributes || info.stores_indexed_attributes ||
info.uses_render_area) {
return true;
}
for (const IR::Block* const block : program.blocks) {
for (const IR::Inst& inst : block->Instructions()) {
switch (inst.GetOpcode()) {
case IR::Opcode::GetAttributeIndexed:
case IR::Opcode::SetAttributeIndexed:
case IR::Opcode::GetPatch:
case IR::Opcode::SetPatch:
return true;
default:
break;
}
}
}
return false;
}
void FinishComputeProgram(IR::Program& program) {
program.stage = Stage::Compute;
program.workgroup_size = {VTG_LOCAL_SIZE, 1, 1};
program.info.uses_workgroup_id = true;
program.info.uses_local_invocation_id = true;
program.info.used_storage_buffer_types |= IR::Type::U32;
program.info.loads = {};
program.info.stores = {};
}
/// Index of the invocation across the whole dispatch
IR::U32 GlobalInvocation(Emitter& ir) {
return ir.Add(ir.IMul(ir.WorkgroupIdX(), ir.Imm32(VTG_LOCAL_SIZE)), ir.LocalInvocationIdX());
}
IR::U32 PrimitiveVertex(Emitter& ir, VtgTopology topology, const IR::U32& primitive, u32 vertex,
const IR::U32& vertex_count) {
const auto offset{[&](const IR::U32& base, u32 value) {
return value == 0 ? base : ir.Add(base, ir.Imm32(value));
}};
const IR::U32 odd{ir.BitwiseAnd(primitive, ir.Imm32(1U))};
const IR::U32 half{ir.ShiftRightLogical(primitive, ir.Imm32(1U))};
switch (topology) {
case VtgTopology::Points:
return primitive;
case VtgTopology::Lines:
return offset(ir.Mul(primitive, 2), vertex);
case VtgTopology::LineLoop:
return ir.Mod(offset(primitive, vertex), ir.UMax(vertex_count, ir.Imm32(1U)));
case VtgTopology::LineStrip:
case VtgTopology::LineStripAdjacency:
return offset(primitive, vertex);
case VtgTopology::Triangles:
return offset(ir.Mul(primitive, 3), vertex);
case VtgTopology::TriangleStrip:
switch (vertex) {
case 0:
return primitive;
case 1:
return ir.Add(offset(primitive, 1), odd);
default:
return ir.Sub(offset(primitive, 2), odd);
}
case VtgTopology::TriangleFan:
return vertex == 0 ? ir.Imm32(0U) : offset(primitive, vertex);
case VtgTopology::Quads: {
const IR::U32 base{ir.Mul(half, 4)};
return vertex == 0 ? base : ir.Add(offset(base, vertex), odd);
}
case VtgTopology::QuadStrip: {
const IR::U32 base{ir.Mul(half, 2)};
switch (vertex) {
case 0:
return base;
case 1:
return ir.Add(offset(base, 1), ir.Mul(odd, 2));
default:
return ir.Sub(offset(base, 3), odd);
}
}
case VtgTopology::LinesAdjacency:
return offset(ir.Mul(primitive, 4), vertex);
case VtgTopology::TrianglesAdjacency:
return offset(ir.Mul(primitive, 6), vertex);
case VtgTopology::TriangleStripAdjacency:
return offset(ir.Mul(primitive, 2), vertex);
}
return primitive;
}
} // Anonymous namespace
bool VtgVertexToCompute(IR::Program& program, const VtgVertexInputs& inputs,
const VtgVaryingLayout& layout, VtgVertexBindings& bindings) {
Info& info{program.info};
if (program.stage != Stage::VertexB || IsIndexedOrUnsupported(program)) {
return false;
}
std::array<bool, VtgDrawInfo::NUM_ATTRIBUTES> uses_buffer{};
size_t num_buffers{};
for (size_t index = 0; index < inputs.size(); ++index) {
const VtgVertexAttribute& attribute{inputs[index]};
if (!info.loads.Generic(index) || attribute.type == VtgComponentType::Disabled) {
continue;
}
if (attribute.components == 0 || attribute.buffer >= uses_buffer.size()) {
return false;
}
num_buffers += uses_buffer[attribute.buffer] ? 0 : 1;
uses_buffer[attribute.buffer] = true;
}
if (info.storage_buffers_descriptors.size() + 3 + num_buffers > Info::MAX_SSBOS) {
return false;
}
bindings.base = static_cast<u32>(info.storage_buffers_descriptors.size());
bindings.draw_info = ReserveStorageBuffer(info, false);
bindings.output = ReserveStorageBuffer(info, true);
bindings.index_buffer = ReserveStorageBuffer(info, false);
std::array<u32, VtgDrawInfo::NUM_ATTRIBUTES> buffer_bindings{};
for (size_t buffer = 0; buffer < uses_buffer.size(); ++buffer) {
if (!uses_buffer[buffer]) {
continue;
}
buffer_bindings[buffer] = ReserveStorageBuffer(info, false);
bindings.vertex_buffers.emplace_back(static_cast<u8>(buffer), buffer_bindings[buffer]);
}
IR::Block& entry{*program.blocks.front()};
Emitter ir{entry, entry.begin()};
const auto draw_info{[&](u32 word) { return ir.LoadWord(bindings.draw_info, word); }};
const IR::U32 invocation{GlobalInvocation(ir)};
const IR::U32 vertex_count{ir.UMax(draw_info(VtgDrawInfo::VERTEX_COUNT), ir.Imm32(1U))};
const IR::U32 last_instance{ir.Sub(draw_info(VtgDrawInfo::INSTANCE_COUNT), ir.Imm32(1U))};
// Invocations past the end of the draw repeat the last instance into unused output
const IR::U32 instance{ir.UMin(ir.IDiv(invocation, vertex_count), last_instance)};
const IR::U32 element{ir.Mod(invocation, vertex_count)};
const IR::U32 base_vertex{draw_info(VtgDrawInfo::BASE_VERTEX)};
const IR::U32 base_instance{draw_info(VtgDrawInfo::BASE_INSTANCE)};
const IR::U32 index_size{draw_info(VtgDrawInfo::INDEX_SIZE)};
const IR::U32 index_address{
ir.Add(draw_info(VtgDrawInfo::INDEX_OFFSET), ir.IMul(element, index_size))};
const IR::U32 index{ir.BitwiseAnd(ir.LoadUnaligned(bindings.index_buffer, index_address, index_size),
draw_info(VtgDrawInfo::INDEX_MASK))};
const IR::U32 vertex_id{ir.Select(ir.IEqual(index_size, ir.Imm32(0U)), element, index)};
const IR::U32 fetch_vertex{ir.Add(vertex_id, base_vertex)};
const IR::U32 output_base{ir.Mul(invocation, layout.stride * WORD_SIZE)};
std::array<std::array<IR::F32, 4>, VtgDrawInfo::NUM_ATTRIBUTES> generics;
for (size_t location = 0; location < inputs.size(); ++location) {
const VtgVertexAttribute& attribute{inputs[location]};
const bool is_read{info.loads.Generic(location) &&
attribute.type != VtgComponentType::Disabled};
IR::U32 address;
if (is_read) {
const u32 words{VtgDrawInfo::ATTRIBUTES +
static_cast<u32>(location) * VtgDrawInfo::ATTRIBUTE_WORDS};
const IR::U32 divisor{
ir.UMax(draw_info(words + VtgDrawInfo::ATTRIBUTE_DIVISOR), ir.Imm32(1U))};
const IR::U32 instanced{ir.Add(base_instance, ir.IDiv(instance, divisor))};
const IR::U1 is_instanced{
ir.INotEqual(draw_info(words + VtgDrawInfo::ATTRIBUTE_INSTANCED), ir.Imm32(0U))};
const IR::U32 source{ir.Select(is_instanced, instanced, fetch_vertex)};
address = ir.Add(draw_info(words + VtgDrawInfo::ATTRIBUTE_OFFSET),
ir.IMul(source, draw_info(words + VtgDrawInfo::ATTRIBUTE_STRIDE)));
}
const u32 binding{buffer_bindings[attribute.buffer % buffer_bindings.size()]};
const bool is_packed{attribute.bits == 0};
const u32 total_bits{static_cast<u32>(attribute.bits) * attribute.components};
// Attributes fitting in a word are read once
const bool single_read{is_read && (is_packed || total_bits <= 32)};
IR::U32 packed;
if (single_read) {
packed = ir.LoadUnaligned(binding, address,
ir.Imm32(is_packed ? WORD_SIZE : total_bits / 8));
}
for (u32 component = 0; component < 4; ++component) {
if (!is_read || !info.loads.Generic(location, component)) {
generics[location][component] = DefaultComponent(ir, attribute, component);
continue;
}
if (component >= attribute.components) {
generics[location][component] = DefaultComponent(ir, attribute, component);
continue;
}
const u32 bits{is_packed ? attribute.packed_bits[component] : attribute.bits};
const u32 bit_offset{is_packed ? attribute.packed_offsets[component]
: component * attribute.bits};
IR::U32 raw;
if (single_read) {
raw = bit_offset == 0
? packed
: IR::U32{ir.ShiftRightLogical(packed, ir.Imm32(bit_offset))};
} else {
raw = ir.LoadUnaligned(binding, ir.Add(address, ir.Imm32(bit_offset / 8)),
ir.Imm32(bits / 8));
}
generics[location][component] = DecodeComponent(ir, attribute, raw, bits);
}
}
for (IR::Block* const block : program.blocks) {
for (IR::Inst& inst : block->Instructions()) {
Emitter at{*block, IteratorTo(inst)};
switch (inst.GetOpcode()) {
case IR::Opcode::Prologue:
case IR::Opcode::Epilogue:
inst.Invalidate();
break;
case IR::Opcode::GetAttribute:
case IR::Opcode::GetAttributeU32: {
const bool is_float{inst.GetOpcode() == IR::Opcode::GetAttribute};
const IR::Attribute attribute{inst.Arg(0).Attribute()};
if (IR::IsGeneric(attribute)) {
const IR::F32 value{generics[IR::GenericAttributeIndex(attribute)]
[IR::GenericAttributeElement(attribute)]};
inst.ReplaceUsesWith(is_float ? IR::Value{value}
: IR::Value{at.BitCast<IR::U32>(value)});
break;
}
IR::U32 value{at.Imm32(0U)};
switch (attribute) {
case IR::Attribute::VertexId:
value = vertex_id;
break;
case IR::Attribute::InstanceId:
value = instance;
break;
case IR::Attribute::BaseVertex:
value = base_vertex;
break;
case IR::Attribute::BaseInstance:
value = base_instance;
break;
default:
break;
}
inst.ReplaceUsesWith(is_float ? IR::Value{at.BitCast<IR::F32>(value)}
: IR::Value{value});
break;
}
case IR::Opcode::SetAttribute: {
const IR::Attribute attribute{inst.Arg(0).Attribute()};
if (const std::optional<u32> word{layout.Offset(attribute)}) {
at.WriteStorage(bindings.output,
at.Add(output_base, at.Imm32(*word * WORD_SIZE)),
at.BitCast<IR::U32>(IR::F32{inst.Arg(1)}));
}
inst.Invalidate();
break;
}
default:
break;
}
}
}
FinishComputeProgram(program);
return true;
}
bool VtgGeometryToCompute(IR::Program& program, VtgTopology topology, u32 input_vertices,
const VtgVaryingLayout& input_layout,
const VtgVaryingLayout& output_layout, VtgGeometryBindings& bindings) {
Info& info{program.info};
if (program.stage != Stage::Geometry || IsIndexedOrUnsupported(program)) {
return false;
}
if (program.invocations > 1 || program.output_vertices == 0 ||
input_vertices > MAX_INPUT_VERTICES || input_vertices == 0) {
return false;
}
if (info.storage_buffers_descriptors.size() + 4 > Info::MAX_SSBOS) {
return false;
}
bindings.base = static_cast<u32>(info.storage_buffers_descriptors.size());
bindings.draw_info = ReserveStorageBuffer(info, false);
bindings.input = ReserveStorageBuffer(info, false);
bindings.output = ReserveStorageBuffer(info, true);
bindings.index_buffer = ReserveStorageBuffer(info, true);
bindings.vertices_per_primitive = program.output_vertices;
// Every vertex may end a strip and the last index always does
bindings.indices_per_primitive = program.output_vertices * 2 + 1;
bindings.output_topology = program.output_topology;
// Emitted vertices are staged in local memory after the memory of the guest
const u32 local_base{Common::DivCeil(program.local_memory_size, WORD_SIZE)};
const u32 vertex_counter{local_base};
const u32 index_counter{local_base + 1};
const u32 staging{local_base + 2};
program.local_memory_size = (staging + output_layout.stride) * WORD_SIZE;
info.uses_local_memory = true;
const u32 max_vertex{bindings.vertices_per_primitive - 1};
const u32 max_index{bindings.indices_per_primitive - 2};
const bool emits_points{program.output_topology == OutputTopology::PointList};
IR::Block& entry{*program.blocks.front()};
Emitter ir{entry, entry.begin()};
const auto draw_info{[&](u32 word) { return ir.LoadWord(bindings.draw_info, word); }};
const IR::U32 invocation{GlobalInvocation(ir)};
const IR::U32 primitive_count{
ir.UMax(draw_info(VtgDrawInfo::PRIMITIVE_COUNT), ir.Imm32(1U))};
const IR::U32 last_instance{ir.Sub(draw_info(VtgDrawInfo::INSTANCE_COUNT), ir.Imm32(1U))};
// Invocations past the end of the draw repeat the last instance into unused output
const IR::U32 instance{ir.UMin(ir.IDiv(invocation, primitive_count), last_instance)};
const IR::U32 primitive{ir.Mod(invocation, primitive_count)};
const IR::U32 vertex_count{draw_info(VtgDrawInfo::VERTEX_COUNT)};
const IR::U32 instance_base{ir.IMul(instance, vertex_count)};
std::array<IR::U32, MAX_INPUT_VERTICES> input_bases;
for (u32 vertex = 0; vertex < input_vertices; ++vertex) {
const IR::U32 source{PrimitiveVertex(ir, topology, primitive, vertex, vertex_count)};
input_bases[vertex] =
ir.Mul(ir.Add(instance_base, source), input_layout.stride * WORD_SIZE);
}
const IR::U32 vertex_base{ir.Mul(invocation, bindings.vertices_per_primitive)};
const IR::U32 index_base{ir.Mul(invocation, bindings.indices_per_primitive)};
ir.WriteLocal(ir.Imm32(vertex_counter), ir.Imm32(0U));
ir.WriteLocal(ir.Imm32(index_counter), ir.Imm32(0U));
const auto input_base{[&](Emitter& at, const IR::Value& vertex) {
if (vertex.IsImmediate()) {
return input_bases[std::min(vertex.U32(), input_vertices - 1)];
}
IR::U32 base{input_bases[0]};
for (u32 index = 1; index < input_vertices; ++index) {
base = IR::U32{at.Select(at.IEqual(IR::U32{vertex}, at.Imm32(index)),
input_bases[index], base)};
}
return base;
}};
const auto index_address{[&](Emitter& at, const IR::U32& count) {
return at.Mul(at.Add(index_base, at.UMin(count, at.Imm32(max_index))), WORD_SIZE);
}};
for (IR::Block* const block : program.blocks) {
for (IR::Inst& inst : block->Instructions()) {
Emitter at{*block, IteratorTo(inst)};
switch (inst.GetOpcode()) {
case IR::Opcode::Prologue:
case IR::Opcode::Epilogue:
inst.Invalidate();
break;
case IR::Opcode::InvocationId:
inst.ReplaceUsesWith(IR::Value{0U});
break;
case IR::Opcode::InvocationInfo:
// Same value the geometry stage gets, the vertices of its input primitive
inst.ReplaceUsesWith(IR::Value{input_vertices << 16});
break;
case IR::Opcode::GetAttribute:
case IR::Opcode::GetAttributeU32: {
const bool is_float{inst.GetOpcode() == IR::Opcode::GetAttribute};
const IR::Attribute attribute{inst.Arg(0).Attribute()};
IR::U32 value;
if (attribute == IR::Attribute::PrimitiveId) {
value = primitive;
} else if (const std::optional<u32> word{input_layout.Offset(attribute)}) {
const IR::U32 base{input_base(at, inst.Arg(1))};
value = at.LoadStorage(bindings.input,
at.Add(base, at.Imm32(*word * WORD_SIZE)));
} else {
value = at.BitCast<IR::U32>(MissingVarying(at, attribute));
}
inst.ReplaceUsesWith(is_float ? IR::Value{at.BitCast<IR::F32>(value)}
: IR::Value{value});
break;
}
case IR::Opcode::SetAttribute: {
const IR::Attribute attribute{inst.Arg(0).Attribute()};
if (const std::optional<u32> word{output_layout.Offset(attribute)}) {
at.WriteLocal(at.Imm32(staging + *word),
at.BitCast<IR::U32>(IR::F32{inst.Arg(1)}));
}
inst.Invalidate();
break;
}
case IR::Opcode::EmitVertex: {
const IR::U32 vertices{at.LoadLocal(at.Imm32(vertex_counter))};
const IR::U32 indices{at.LoadLocal(at.Imm32(index_counter))};
const IR::U32 vertex{
at.Add(vertex_base, at.UMin(vertices, at.Imm32(max_vertex)))};
const IR::U32 address{at.Mul(vertex, output_layout.stride * WORD_SIZE)};
for (u32 word = 0; word < output_layout.stride; ++word) {
at.WriteStorage(bindings.output, at.Add(address, at.Imm32(word * WORD_SIZE)),
at.LoadLocal(at.Imm32(staging + word)));
}
at.WriteStorage(bindings.index_buffer, index_address(at, indices), vertex);
at.WriteLocal(at.Imm32(vertex_counter), at.Add(vertices, at.Imm32(1U)));
at.WriteLocal(at.Imm32(index_counter), at.Add(indices, at.Imm32(1U)));
inst.Invalidate();
break;
}
case IR::Opcode::EndPrimitive: {
if (!emits_points) {
const IR::U32 indices{at.LoadLocal(at.Imm32(index_counter))};
at.WriteStorage(bindings.index_buffer, index_address(at, indices),
at.Imm32(VTG_RESTART_INDEX));
at.WriteLocal(at.Imm32(index_counter), at.Add(indices, at.Imm32(1U)));
}
inst.Invalidate();
break;
}
default:
break;
}
}
}
FinishComputeProgram(program);
return true;
}
} // namespace Shader::Optimization

3
src/yuzu_shader_recompiler/profile.h
View File

@ -93,6 +93,9 @@ struct Profile {
u32 max_user_clip_distances{};
/// Largest compute shared memory the host accepts, in bytes. Zero means no limit.
u32 max_compute_shared_memory_size{};
bool SupportsUniformBufferDescriptorAliasing() const {
return support_descriptor_aliasing && !disable_uniform_buffer_descriptor_aliasing;
}

5
src/yuzu_shader_recompiler/runtime_info.h
View File

@ -9,6 +9,7 @@
#include <vector>
#include "yuzu_common/common_types.h"
#include "yuzu_shader_recompiler/frontend/ir/attribute.h"
#include "yuzu_shader_recompiler/varying_state.h"
namespace Shader {
@ -90,6 +91,10 @@ struct RuntimeInfo {
TessSpacing tess_spacing{};
bool tess_clockwise{};
/// Generic output copied into the layer at the end of the vertex stage, standing in for an
/// eliminated geometry shader
std::optional<IR::Attribute> forwarded_layer_attribute;
InputTopology input_topology{};
std::optional<float> fixed_state_point_size;

263
src/yuzu_shader_recompiler/vtg_as_compute.h Normal file
View File

@ -0,0 +1,263 @@
// SPDX-FileCopyrightText: Copyright 2026 NxEmu Project
// SPDX-License-Identifier: GPL-2.0-or-later
#pragma once
#include <array>
#include <optional>
#include <boost/container/static_vector.hpp>
#include "yuzu_common/common_types.h"
#include "yuzu_shader_recompiler/frontend/ir/attribute.h"
#include "yuzu_shader_recompiler/program_header.h"
#include "yuzu_shader_recompiler/varying_state.h"
// On hosts without geometry shaders, the vertex and geometry stages of a draw run as compute
// shaders. The vertex stage pulls its attributes from the vertex buffers and writes its outputs to
// a storage buffer, the geometry stage reads them back and writes the vertices it emits with a
// strip index buffer, and a generated vertex shader feeds the result to the rasterizer.
namespace Shader {
/// Invocations per workgroup of the stages running as compute
inline constexpr u32 VTG_LOCAL_SIZE = 32;
/// Index marking the end of a strip in the generated index buffer
inline constexpr u32 VTG_RESTART_INDEX = 0xffffffff;
/// Words of the buffer describing the draw to the stages running as compute
namespace VtgDrawInfo {
inline constexpr u32 VERTEX_COUNT = 0;
inline constexpr u32 INSTANCE_COUNT = 1;
inline constexpr u32 BASE_VERTEX = 2;
inline constexpr u32 BASE_INSTANCE = 3;
/// Bytes per index, zero on non indexed draws
inline constexpr u32 INDEX_SIZE = 4;
/// Byte offset of the first index inside the bound index buffer
inline constexpr u32 INDEX_OFFSET = 5;
inline constexpr u32 PRIMITIVE_COUNT = 6;
inline constexpr u32 INDEX_MASK = 7;
/// First word of the four words describing each vertex attribute
inline constexpr u32 ATTRIBUTES = 8;
inline constexpr u32 ATTRIBUTE_OFFSET = 0;
inline constexpr u32 ATTRIBUTE_STRIDE = 1;
inline constexpr u32 ATTRIBUTE_INSTANCED = 2;
inline constexpr u32 ATTRIBUTE_DIVISOR = 3;
inline constexpr u32 ATTRIBUTE_WORDS = 4;
inline constexpr u32 NUM_ATTRIBUTES = 32;
inline constexpr u32 NUM_WORDS = ATTRIBUTES + NUM_ATTRIBUTES * ATTRIBUTE_WORDS;
} // namespace VtgDrawInfo
enum class VtgComponentType : u8 {
Disabled,
Float,
UNorm,
SNorm,
UInt,
SInt,
UScaled,
SScaled,
};
struct VtgVertexAttribute {
VtgComponentType type{VtgComponentType::Disabled};
/// Vertex buffer the attribute is read from
u8 buffer{};
u8 components{};
/// Bits per component, zero when the components are packed in a single word
u8 bits{};
/// Bit position and size of each component of a packed format
std::array<u8, 4> packed_offsets{};
std::array<u8, 4> packed_bits{};
};
using VtgVertexInputs = std::array<VtgVertexAttribute, VtgDrawInfo::NUM_ATTRIBUTES>;
/// Primitives assembled from the vertices produced by the vertex stage
enum class VtgTopology : u8 {
Points,
Lines,
LineLoop,
LineStrip,
Triangles,
TriangleStrip,
TriangleFan,
Quads,
QuadStrip,
LinesAdjacency,
LineStripAdjacency,
TrianglesAdjacency,
TriangleStripAdjacency,
};
/// Number of primitives the geometry stage receives from a number of vertices
[[nodiscard]] constexpr u32 VtgPrimitiveCount(VtgTopology topology, u32 count) noexcept {
switch (topology) {
case VtgTopology::Points:
return count;
case VtgTopology::Lines:
return count / 2;
case VtgTopology::LineLoop:
return count >= 2 ? count : 0;
case VtgTopology::LineStrip:
return count >= 2 ? count - 1 : 0;
case VtgTopology::Triangles:
return count / 3;
case VtgTopology::TriangleStrip:
case VtgTopology::TriangleFan:
return count >= 3 ? count - 2 : 0;
case VtgTopology::Quads:
return (count / 4) * 2;
case VtgTopology::QuadStrip:
return count >= 4 ? ((count - 2) / 2) * 2 : 0;
case VtgTopology::LinesAdjacency:
return count / 4;
case VtgTopology::LineStripAdjacency:
return count >= 4 ? count - 3 : 0;
case VtgTopology::TrianglesAdjacency:
return count / 6;
case VtgTopology::TriangleStripAdjacency:
return count >= 6 ? (count - 4) / 2 : 0;
}
return 0;
}
/// Vertex produced by the vertex stage used as a vertex of a primitive
[[nodiscard]] constexpr u32 VtgPrimitiveVertex(VtgTopology topology, u32 primitive, u32 vertex,
u32 count) noexcept {
switch (topology) {
case VtgTopology::Points:
return primitive;
case VtgTopology::Lines:
return primitive * 2 + vertex;
case VtgTopology::LineLoop:
return count != 0 ? (primitive + vertex) % count : 0;
case VtgTopology::LineStrip:
case VtgTopology::LineStripAdjacency:
return primitive + vertex;
case VtgTopology::Triangles:
return primitive * 3 + vertex;
case VtgTopology::TriangleStrip: {
// Odd triangles swap their last two vertices to keep the winding
const u32 odd{primitive & 1};
const u32 offset{vertex == 0 ? 0U : (vertex == 1 ? 1U + odd : 2U - odd)};
return primitive + offset;
}
case VtgTopology::TriangleFan:
return vertex == 0 ? 0 : primitive + vertex;
case VtgTopology::Quads: {
// Each quad is the triangles (0, 1, 2) and (0, 2, 3)
const u32 second{primitive & 1};
const u32 base{(primitive / 2) * 4};
return vertex == 0 ? base : base + vertex + second;
}
case VtgTopology::QuadStrip: {
// Each quad is the triangles (0, 1, 3) and (0, 3, 2)
const u32 second{primitive & 1};
const u32 base{(primitive / 2) * 2};
const u32 offset{vertex == 0 ? 0U : (vertex == 1 ? 1U + second * 2U : 3U - second)};
return base + offset;
}
case VtgTopology::LinesAdjacency:
return primitive * 4 + vertex;
case VtgTopology::TrianglesAdjacency:
return primitive * 6 + vertex;
case VtgTopology::TriangleStripAdjacency:
return primitive * 2 + vertex;
}
return 0;
}
/// Layout in words of a vertex in the buffers shared between stages
struct VtgVaryingLayout {
static constexpr u32 POSITION_WORDS = 4;
static constexpr u32 GENERIC_WORDS = 4;
static constexpr u32 CLIP_DISTANCE_WORDS = 8;
static constexpr u32 NUM_GENERICS = 32;
std::array<s32, NUM_GENERICS> generics{};
s32 point_size{-1};
s32 clip_distances{-1};
s32 layer{-1};
s32 viewport_index{-1};
u32 stride{};
explicit VtgVaryingLayout() {
generics.fill(-1);
}
explicit VtgVaryingLayout(const VaryingState& stores) : VtgVaryingLayout() {
stride = POSITION_WORDS;
for (size_t index = 0; index < NUM_GENERICS; ++index) {
if (stores.Generic(index)) {
generics[index] = static_cast<s32>(stride);
stride += GENERIC_WORDS;
}
}
const auto reserve{[this](bool used, u32 words) {
const s32 offset{used ? static_cast<s32>(stride) : -1};
stride += used ? words : 0;
return offset;
}};
point_size = reserve(stores[IR::Attribute::PointSize], 1);
clip_distances = reserve(stores.ClipDistances(), CLIP_DISTANCE_WORDS);
layer = reserve(stores[IR::Attribute::Layer], 1);
viewport_index = reserve(stores[IR::Attribute::ViewportIndex], 1);
}
/// Word offset of an attribute inside a vertex, if the layout has it
[[nodiscard]] std::optional<u32> Offset(IR::Attribute attribute) const {
const auto at{[](s32 base, u32 element) -> std::optional<u32> {
return base < 0 ? std::nullopt : std::optional<u32>{static_cast<u32>(base) + element};
}};
if (attribute >= IR::Attribute::PositionX && attribute <= IR::Attribute::PositionW) {
return static_cast<u32>(attribute) - static_cast<u32>(IR::Attribute::PositionX);
}
if (IR::IsGeneric(attribute)) {
return at(generics[IR::GenericAttributeIndex(attribute)],
IR::GenericAttributeElement(attribute));
}
if (attribute >= IR::Attribute::ClipDistance0 && attribute <= IR::Attribute::ClipDistance7) {
return at(clip_distances, static_cast<u32>(attribute) -
static_cast<u32>(IR::Attribute::ClipDistance0));
}
switch (attribute) {
case IR::Attribute::PointSize:
return at(point_size, 0);
case IR::Attribute::Layer:
return at(layer, 0);
case IR::Attribute::ViewportIndex:
return at(viewport_index, 0);
default:
return std::nullopt;
}
}
};
/// Storage buffers added to a vertex stage running as compute
struct VtgVertexBindings {
/// Index of the first added storage buffer descriptor
u32 base{};
u32 draw_info{};
u32 output{};
u32 index_buffer{};
/// Vertex buffers read by the stage and the storage buffer each one is bound to
boost::container::static_vector<std::pair<u8, u32>, VtgDrawInfo::NUM_ATTRIBUTES> vertex_buffers;
};
/// Storage buffers added to a geometry stage running as compute
struct VtgGeometryBindings {
/// Index of the first added storage buffer descriptor
u32 base{};
u32 draw_info{};
u32 input{};
u32 output{};
u32 index_buffer{};
/// Vertices and indices reserved for each input primitive
u32 vertices_per_primitive{};
u32 indices_per_primitive{};
OutputTopology output_topology{};
};
} // namespace Shader

4
src/yuzu_shader_recompiler/yuzu_shader_recompiler.vcxproj
View File

@ -85,6 +85,7 @@
<ClInclude Include="shader_info.h" />
<ClInclude Include="stage.h" />
<ClInclude Include="varying_state.h" />
<ClInclude Include="vtg_as_compute.h" />
</ItemGroup>
<ItemGroup>
<ClCompile Include="backend\glasm\emit_glasm.cpp" />
@ -269,10 +270,13 @@
<ClCompile Include="ir_opt\lower_int64_to_int32.cpp" />
<ClCompile Include="ir_opt\position_pass.cpp" />
<ClCompile Include="ir_opt\rescaling_pass.cpp" />
<ClCompile Include="ir_opt\geometry_layer_forwarding_pass.cpp" />
<ClCompile Include="ir_opt\shared_memory_packing_pass.cpp" />
<ClCompile Include="ir_opt\ssa_rewrite_pass.cpp" />
<ClCompile Include="ir_opt\texture_pass.cpp" />
<ClCompile Include="ir_opt\vendor_workaround_pass.cpp" />
<ClCompile Include="ir_opt\verification_pass.cpp" />
<ClCompile Include="ir_opt\vtg_to_compute_pass.cpp" />
</ItemGroup>
<Import Project="$(VCTargetsPath)\Microsoft.Cpp.targets" />
</Project>

12
src/yuzu_shader_recompiler/yuzu_shader_recompiler.vcxproj.filters
View File

@ -104,6 +104,9 @@
<ClInclude Include="varying_state.h">
<Filter>Header Files</Filter>
</ClInclude>
<ClInclude Include="vtg_as_compute.h">
<Filter>Header Files</Filter>
</ClInclude>
<ClInclude Include="backend\bindings.h">
<Filter>Header Files\backend</Filter>
</ClInclude>
@ -772,6 +775,12 @@
<ClCompile Include="ir_opt\rescaling_pass.cpp">
<Filter>Source Files\ir_opt</Filter>
</ClCompile>
<ClCompile Include="ir_opt\geometry_layer_forwarding_pass.cpp">
<Filter>Source Files\ir_opt</Filter>
</ClCompile>
<ClCompile Include="ir_opt\shared_memory_packing_pass.cpp">
<Filter>Source Files\ir_opt</Filter>
</ClCompile>
<ClCompile Include="ir_opt\ssa_rewrite_pass.cpp">
<Filter>Source Files\ir_opt</Filter>
</ClCompile>
@ -784,5 +793,8 @@
<ClCompile Include="ir_opt\verification_pass.cpp">
<Filter>Source Files\ir_opt</Filter>
</ClCompile>
<ClCompile Include="ir_opt\vtg_to_compute_pass.cpp">
<Filter>Source Files\ir_opt</Filter>
</ClCompile>
</ItemGroup>
</Project>

37
src/yuzu_video_core/buffer_cache/buffer_cache.h
View File

@ -370,6 +370,43 @@ void BufferCache<P>::BindHostStageBuffers(size_t stage) {
BindHostGraphicsTextureBuffers(stage);
}
template <class P>
void BufferCache<P>::BindHostStageUniformAndStorageBuffers(size_t stage) {
BindHostGraphicsUniformBuffers(stage);
BindHostGraphicsStorageBuffers(stage);
}
template <class P>
void BufferCache<P>::BindHostStageTextureBuffers(size_t stage) {
BindHostGraphicsTextureBuffers(stage);
}
template <class P>
std::tuple<typename P::Buffer*, u32, u32> BufferCache<P>::VertexBufferBinding(u32 index) {
const Binding& binding = channel_state->vertex_buffers[index];
if (binding.buffer_id == NULL_BUFFER_ID || binding.size == 0) {
return {nullptr, 0, 0};
}
Buffer& buffer = slot_buffers[binding.buffer_id];
TouchBuffer(buffer, binding.buffer_id);
SynchronizeBuffer(buffer, binding.device_addr, binding.size);
const u32 offset = buffer.Offset(binding.device_addr);
buffer.MarkUsage(offset, binding.size);
return {&buffer, offset, binding.size};
}
template <class P>
std::tuple<typename P::Buffer*, u32, u32> BufferCache<P>::IndexBufferBinding() {
const Binding& binding = channel_state->index_buffer;
if (binding.buffer_id == NULL_BUFFER_ID || binding.size == 0) {
return {nullptr, 0, 0};
}
Buffer& buffer = slot_buffers[binding.buffer_id];
const u32 offset = buffer.Offset(binding.device_addr);
buffer.MarkUsage(offset, binding.size);
return {&buffer, offset, binding.size};
}
template <class P>
void BufferCache<P>::BindHostComputeBuffers() {
BindHostComputeUniformBuffers();

13
src/yuzu_video_core/buffer_cache/buffer_cache_base.h
View File

@ -5,6 +5,7 @@
#include <algorithm>
#include <array>
#include <tuple>
#include <functional>
#include <memory>
#include <mutex>
@ -218,6 +219,18 @@ public:
void BindHostStageBuffers(size_t stage);
/// Binds the uniform and storage buffers of a stage, leaving its texture buffers unbound
void BindHostStageUniformAndStorageBuffers(size_t stage);
/// Binds the texture buffers of a stage
void BindHostStageTextureBuffers(size_t stage);
/// Returns the host buffer with the offset and size of a vertex stream of the draw
[[nodiscard]] std::tuple<Buffer*, u32, u32> VertexBufferBinding(u32 index);
/// Returns the host buffer with the offset and size of the index buffer of the draw
[[nodiscard]] std::tuple<Buffer*, u32, u32> IndexBufferBinding();
void BindHostComputeBuffers();
void SetUniformBuffersState(const std::array<u32, NUM_STAGES>& mask,

80
src/yuzu_video_core/renderer_vulkan/vk_compute_pipeline.cpp
View File

@ -41,46 +41,51 @@ ComputePipeline::ComputePipeline(const Device& device_, vk::PipelineCache& pipel
uniform_buffer_sizes.begin());
auto func{[this, &descriptor_pool, shader_notify, pipeline_statistics] {
DescriptorLayoutBuilder builder{device};
builder.Add(info, VK_SHADER_STAGE_COMPUTE_BIT);
try {
DescriptorLayoutBuilder builder{device};
builder.Add(info, VK_SHADER_STAGE_COMPUTE_BIT);
descriptor_set_layout = builder.CreateDescriptorSetLayout(false);
pipeline_layout = builder.CreatePipelineLayout(*descriptor_set_layout);
descriptor_update_template =
builder.CreateTemplate(*descriptor_set_layout, *pipeline_layout, false);
descriptor_allocator = descriptor_pool.Allocator(*descriptor_set_layout, info);
const VkPipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroup_size_ci{
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_REQUIRED_SUBGROUP_SIZE_CREATE_INFO_EXT,
.pNext = nullptr,
.requiredSubgroupSize = GuestWarpSize,
};
VkPipelineCreateFlags flags{};
if (device.IsKhrPipelineExecutablePropertiesEnabled()) {
flags |= VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR;
}
pipeline = device.GetLogical().CreateComputePipeline(
{
.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
descriptor_set_layout = builder.CreateDescriptorSetLayout(false);
pipeline_layout = builder.CreatePipelineLayout(*descriptor_set_layout);
descriptor_update_template =
builder.CreateTemplate(*descriptor_set_layout, *pipeline_layout, false);
descriptor_allocator = descriptor_pool.Allocator(*descriptor_set_layout, info);
const VkPipelineShaderStageRequiredSubgroupSizeCreateInfoEXT subgroup_size_ci{
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_REQUIRED_SUBGROUP_SIZE_CREATE_INFO_EXT,
.pNext = nullptr,
.flags = flags,
.stage{
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
.pNext =
device.CanRequestComputeSubgroupSize() ? &subgroup_size_ci : nullptr,
.flags = 0,
.stage = VK_SHADER_STAGE_COMPUTE_BIT,
.module = *spv_module,
.pName = "main",
.pSpecializationInfo = nullptr,
.requiredSubgroupSize = GuestWarpSize,
};
VkPipelineCreateFlags flags{};
if (device.IsKhrPipelineExecutablePropertiesEnabled()) {
flags |= VK_PIPELINE_CREATE_CAPTURE_STATISTICS_BIT_KHR;
}
pipeline = device.GetLogical().CreateComputePipeline(
{
.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
.pNext = nullptr,
.flags = flags,
.stage{
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
.pNext =
device.CanRequestComputeSubgroupSize() ? &subgroup_size_ci : nullptr,
.flags = 0,
.stage = VK_SHADER_STAGE_COMPUTE_BIT,
.module = *spv_module,
.pName = "main",
.pSpecializationInfo = nullptr,
},
.layout = *pipeline_layout,
.basePipelineHandle = 0,
.basePipelineIndex = 0,
},
.layout = *pipeline_layout,
.basePipelineHandle = 0,
.basePipelineIndex = 0,
},
*pipeline_cache);
*pipeline_cache);
if (pipeline_statistics) {
pipeline_statistics->Collect(*pipeline);
if (pipeline_statistics) {
pipeline_statistics->Collect(*pipeline);
}
} catch (const vk::Exception& exception) {
LOG_ERROR(Render_Vulkan, "Compute pipeline build failed: {}", exception.what());
build_failed = true;
}
std::scoped_lock lock{build_mutex};
is_built = true;
@ -208,6 +213,9 @@ void ComputePipeline::Configure(Tegra::Engines::KeplerCompute& kepler_compute,
const bool is_rescaling = !info.texture_descriptors.empty() || !info.image_descriptors.empty();
scheduler.Record([this, descriptor_data, is_rescaling,
rescaling_data = rescaling.Data()](vk::CommandBuffer cmdbuf) {
if (build_failed) {
return;
}
cmdbuf.BindPipeline(VK_PIPELINE_BIND_POINT_COMPUTE, *pipeline);
if (!descriptor_set_layout) {
return;

6
src/yuzu_video_core/renderer_vulkan/vk_compute_pipeline.h
View File

@ -45,6 +45,11 @@ public:
void Configure(Tegra::Engines::KeplerCompute& kepler_compute, Tegra::MemoryManager& gpu_memory,
Scheduler& scheduler, BufferCache& buffer_cache, TextureCache& texture_cache);
/// Returns true when the host rejected the pipeline. Only valid once the build has finished.
[[nodiscard]] bool HasBuildFailed() const noexcept {
return build_failed.load(std::memory_order::relaxed);
}
private:
const Device& device;
vk::PipelineCache& pipeline_cache;
@ -63,6 +68,7 @@ private:
std::condition_variable build_condvar;
std::mutex build_mutex;
std::atomic_bool is_built{false};
std::atomic_bool build_failed{false};
};
} // namespace Vulkan

497
src/yuzu_video_core/renderer_vulkan/vk_graphics_pipeline.cpp
View File

@ -9,6 +9,7 @@
#include "yuzu_video_core/renderer_vulkan/pipeline_helper.h"
#include "yuzu_common/alignment.h"
#include "yuzu_common/bit_field.h"
#include "yuzu_video_core/renderer_vulkan/maxwell_to_vk.h"
#include "yuzu_video_core/renderer_vulkan/pipeline_statistics.h"
@ -16,6 +17,7 @@
#include "yuzu_video_core/renderer_vulkan/vk_graphics_pipeline.h"
#include "yuzu_video_core/renderer_vulkan/vk_render_pass_cache.h"
#include "yuzu_video_core/renderer_vulkan/vk_scheduler.h"
#include "yuzu_video_core/renderer_vulkan/vk_staging_buffer_pool.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/shader_notify.h"
@ -271,6 +273,78 @@ ConfigureFuncPtr ConfigureFunc(const std::array<vk::ShaderModule, NUM_STAGES>& m
}
} // Anonymous namespace
VtgScratchAllocator::VtgScratchAllocator(const Device& device, Scheduler& scheduler_,
StagingBufferPool& staging_pool_)
: scheduler{scheduler_}, staging_pool{staging_pool_},
alignment{std::max<VkDeviceSize>(device.GetStorageBufferAlignment(), sizeof(u32))} {}
std::pair<VkBuffer, VkDeviceSize> VtgScratchAllocator::Allocate(VkDeviceSize size) {
static constexpr VkDeviceSize CHUNK_SIZE = 4ULL << 20;
size = Common::AlignUp(size, alignment);
const u64 current_tick{scheduler.CurrentTick()};
if (buffer == VK_NULL_HANDLE || tick != current_tick || used + size > capacity) {
// A chunk is only safe to write until the tick that requested it completes
capacity = std::max(CHUNK_SIZE, size);
const StagingBufferRef ref{staging_pool.Request(capacity, MemoryUsage::DeviceLocal)};
buffer = ref.buffer;
base = ref.offset;
used = 0;
tick = current_tick;
}
const VkDeviceSize offset{base + used};
used += size;
return {buffer, offset};
}
struct GraphicsPipeline::VtgState {
struct Stage {
/// Shader information including the storage buffers added to the stage
Shader::Info info;
/// Shader information of the resources bound by the guest
Shader::Info guest_info;
vk::ShaderModule module;
vk::DescriptorSetLayout descriptor_set_layout;
DescriptorAllocator descriptor_allocator;
vk::PipelineLayout pipeline_layout;
vk::DescriptorUpdateTemplate descriptor_update_template;
vk::Pipeline pipeline;
};
struct Buffer {
VkBuffer buffer;
VkDeviceSize offset;
VkDeviceSize size;
};
std::array<Stage, VtgPipelineInfo::NUM_COMPUTE_STAGES> stages;
Shader::VtgVertexBindings vertex_bindings;
Shader::VtgGeometryBindings geometry_bindings;
Shader::VtgTopology topology{};
u32 vertex_stride{};
u32 geometry_stride{};
VtgScratchAllocator* scratch{};
StagingBufferPool* staging_pool{};
std::array<u32, NUM_STAGES> uniform_masks{};
VideoCommon::UniformBufferSizes uniform_sizes{};
u32 instance_count{};
VkBuffer index_buffer{};
VkDeviceSize index_offset{};
u32 index_count{};
};
GraphicsPipeline::~GraphicsPipeline() = default;
void GraphicsPipeline::SetVtgInstanceCount(u32 instance_count) noexcept {
if (vtg) {
vtg->instance_count = instance_count;
}
}
std::tuple<VkBuffer, VkDeviceSize, u32> GraphicsPipeline::VtgDraw() const noexcept {
return {vtg->index_buffer, vtg->index_offset, vtg->index_count};
}
GraphicsPipeline::GraphicsPipeline(
Scheduler& scheduler_, BufferCache& buffer_cache_, TextureCache& texture_cache_,
vk::PipelineCache& pipeline_cache_, VideoCore::ShaderNotify* shader_notify,
@ -278,7 +352,8 @@ GraphicsPipeline::GraphicsPipeline(
GuestDescriptorQueue& guest_descriptor_queue_, Common::ThreadWorker* worker_thread,
PipelineStatistics* pipeline_statistics, RenderPassCache& render_pass_cache,
const GraphicsPipelineCacheKey& key_, std::array<vk::ShaderModule, NUM_STAGES> stages,
const std::array<const Shader::Info*, NUM_STAGES>& infos, bool rasterizes_lines_)
const std::array<const Shader::Info*, NUM_STAGES>& infos, bool rasterizes_lines_,
std::unique_ptr<VtgPipelineInfo> vtg_info)
: key{key_}, device{device_}, texture_cache{texture_cache_}, buffer_cache{buffer_cache_},
pipeline_cache(pipeline_cache_), scheduler{scheduler_},
guest_descriptor_queue{guest_descriptor_queue_}, spv_modules{std::move(stages)},
@ -296,6 +371,32 @@ GraphicsPipeline::GraphicsPipeline(
std::ranges::copy(info->constant_buffer_used_sizes, uniform_buffer_sizes[stage].begin());
num_textures += Shader::NumDescriptors(info->texture_descriptors);
}
if (vtg_info) {
static constexpr std::array<size_t, VtgPipelineInfo::NUM_COMPUTE_STAGES> compute_stages{
0, 3};
const std::array<u32, VtgPipelineInfo::NUM_COMPUTE_STAGES> guest_storage_buffers{
vtg_info->vertex_bindings.base, vtg_info->geometry_bindings.base};
vtg = std::make_unique<VtgState>();
vtg->uniform_masks = enabled_uniform_buffer_masks;
vtg->uniform_sizes = uniform_buffer_sizes;
for (size_t index = 0; index < vtg->stages.size(); ++index) {
VtgState::Stage& stage{vtg->stages[index]};
stage.info = vtg_info->infos[index];
stage.guest_info = stage.info;
stage.guest_info.storage_buffers_descriptors.resize(guest_storage_buffers[index]);
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,
vtg->uniform_sizes[compute_stages[index]].begin());
}
vtg->vertex_bindings = vtg_info->vertex_bindings;
vtg->geometry_bindings = vtg_info->geometry_bindings;
vtg->topology = vtg_info->topology;
vtg->vertex_stride = vtg_info->vertex_stride;
vtg->geometry_stride = vtg_info->geometry_stride;
vtg->scratch = vtg_info->scratch;
vtg->staging_pool = vtg_info->staging_pool;
}
// 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]};
@ -329,6 +430,38 @@ GraphicsPipeline::GraphicsPipeline(
descriptor_update_template =
builder.CreateTemplate(set_layout, *pipeline_layout, uses_push_descriptor);
if (vtg) {
for (VtgState::Stage& stage : vtg->stages) {
DescriptorLayoutBuilder stage_builder{device};
stage_builder.Add(stage.info, VK_SHADER_STAGE_COMPUTE_BIT);
stage.descriptor_set_layout = stage_builder.CreateDescriptorSetLayout(false);
stage.pipeline_layout =
stage_builder.CreatePipelineLayout(*stage.descriptor_set_layout);
stage.descriptor_update_template = stage_builder.CreateTemplate(
*stage.descriptor_set_layout, *stage.pipeline_layout, false);
stage.descriptor_allocator =
descriptor_pool.Allocator(*stage.descriptor_set_layout, stage.info);
stage.pipeline = device.GetLogical().CreateComputePipeline(
{
.sType = VK_STRUCTURE_TYPE_COMPUTE_PIPELINE_CREATE_INFO,
.pNext = nullptr,
.flags = 0,
.stage{
.sType = VK_STRUCTURE_TYPE_PIPELINE_SHADER_STAGE_CREATE_INFO,
.pNext = nullptr,
.flags = 0,
.stage = VK_SHADER_STAGE_COMPUTE_BIT,
.module = *stage.module,
.pName = "main",
.pSpecializationInfo = nullptr,
},
.layout = *stage.pipeline_layout,
.basePipelineHandle = 0,
.basePipelineIndex = 0,
},
*pipeline_cache);
}
}
const VkRenderPass render_pass{render_pass_cache.Get(MakeRenderPassKey(key.state))};
Validate();
MakePipeline(render_pass);
@ -356,7 +489,11 @@ GraphicsPipeline::GraphicsPipeline(
} else {
func();
}
configure_func = ConfigureFunc(spv_modules, stage_infos);
if (vtg) {
configure_func = [](GraphicsPipeline* pl, bool is_indexed) { pl->ConfigureVtg(is_indexed); };
} else {
configure_func = ConfigureFunc(spv_modules, stage_infos);
}
}
// A pipeline can be shared by several keys, so the transition is keyed explicitly.
@ -557,6 +694,346 @@ void GraphicsPipeline::ConfigureImpl(bool is_indexed) {
ConfigureDraw(rescaling, render_area);
}
void GraphicsPipeline::ConfigureVtg(bool is_indexed) {
using Shader::VTG_LOCAL_SIZE;
namespace DrawInfo = Shader::VtgDrawInfo;
static constexpr size_t FRAGMENT_STAGE = 4;
static constexpr u32 MAX_INVOCATIONS = 1U << 20;
static constexpr VkDeviceSize MAX_SCRATCH_SIZE = 256ULL << 20;
VtgState& state{*vtg};
state.index_count = 0;
std::array<VideoCommon::ImageViewInOut, MAX_IMAGE_ELEMENTS> views;
std::array<VideoCommon::SamplerId, MAX_IMAGE_ELEMENTS> samplers;
size_t sampler_index{};
size_t view_index{};
texture_cache.SynchronizeGraphicsDescriptors();
buffer_cache.SetUniformBuffersState(state.uniform_masks, &state.uniform_sizes);
// The guest resources of the stages running as compute are bound as the stages they replace
const std::array<std::pair<size_t, const Shader::Info*>, 3> stages{{
{0, &state.stages[VtgPipelineInfo::VERTEX].guest_info},
{3, &state.stages[VtgPipelineInfo::GEOMETRY].guest_info},
{FRAGMENT_STAGE, &stage_infos[FRAGMENT_STAGE]},
}};
const auto& regs{maxwell3d->regs};
const bool via_header_index{regs.sampler_binding == Maxwell::SamplerBinding::ViaHeaderBinding};
for (const auto& [stage, info_pointer] : stages) {
const Shader::Info& info{*info_pointer};
buffer_cache.UnbindGraphicsStorageBuffers(stage);
size_t ssbo_index{};
for (const auto& desc : info.storage_buffers_descriptors) {
ASSERT(desc.count == 1);
buffer_cache.BindGraphicsStorageBuffer(stage, ssbo_index, desc.cbuf_index,
desc.cbuf_offset, desc.is_written);
++ssbo_index;
}
const auto& cbufs{maxwell3d->state.shader_stages[stage].const_buffers};
const auto read_handle{[&](const auto& desc, u32 index) {
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&>) {
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};
return TexturePair(lhs_raw | rhs_raw, via_header_index);
}
}
return TexturePair(gpu_memory->Read<u32>(addr), via_header_index);
}};
const auto add_image{[&](const auto& desc, bool blacklist) {
for (u32 index = 0; index < desc.count; ++index) {
const auto handle{read_handle(desc, index)};
views[view_index++] = {
.index = handle.first,
.blacklist = blacklist,
.id = {},
};
}
}};
for (const auto& desc : info.texture_buffer_descriptors) {
add_image(desc, false);
}
for (const auto& desc : info.image_buffer_descriptors) {
add_image(desc, false);
}
for (const auto& desc : info.texture_descriptors) {
for (u32 index = 0; index < desc.count; ++index) {
const auto handle{read_handle(desc, index)};
views[view_index++] = {handle.first};
samplers[sampler_index++] = texture_cache.GetGraphicsSamplerId(handle.second);
}
}
for (const auto& desc : info.image_descriptors) {
add_image(desc, desc.is_written);
}
}
texture_cache.FillGraphicsImageViews<true>(std::span(views.data(), view_index));
VideoCommon::ImageViewInOut* texture_buffer_it{views.data()};
for (const auto& [stage, info_pointer] : stages) {
const Shader::Info& info{*info_pointer};
size_t index{};
const auto add_buffer{[&](const auto& desc) {
constexpr bool is_image = std::is_same_v<decltype(desc), const ImageBufferDescriptor&>;
for (u32 i = 0; i < desc.count; ++i) {
bool is_written{false};
if constexpr (is_image) {
is_written = desc.is_written;
}
ImageView& image_view{texture_cache.GetImageView(texture_buffer_it->id)};
buffer_cache.BindGraphicsTextureBuffer(stage, index, image_view.GpuAddr(),
image_view.BufferSize(), image_view.format,
is_written, is_image);
++index;
++texture_buffer_it;
}
}};
buffer_cache.UnbindGraphicsTextureBuffers(stage);
for (const auto& desc : info.texture_buffer_descriptors) {
add_buffer(desc);
}
for (const auto& desc : info.image_buffer_descriptors) {
add_buffer(desc);
}
texture_buffer_it += Shader::NumDescriptors(info.texture_descriptors);
texture_buffer_it += Shader::NumDescriptors(info.image_descriptors);
}
buffer_cache.UpdateGraphicsBuffers(is_indexed);
buffer_cache.BindHostGeometryBuffers(is_indexed);
const auto& draw_state{maxwell3d->draw_manager->GetDrawState()};
const u32 vertex_count{is_indexed ? draw_state.index_buffer.count
: draw_state.vertex_buffer.count};
const u32 instance_count{state.instance_count};
const u32 primitive_count{Shader::VtgPrimitiveCount(state.topology, vertex_count)};
const u64 total_vertices{u64{vertex_count} * instance_count};
const u64 total_primitives{u64{primitive_count} * instance_count};
const u32 vertices_per_primitive{state.geometry_bindings.vertices_per_primitive};
const u32 indices_per_primitive{state.geometry_bindings.indices_per_primitive};
const bool fits{total_vertices <= MAX_INVOCATIONS && total_primitives <= MAX_INVOCATIONS &&
(total_vertices * state.vertex_stride +
total_primitives * (u64{vertices_per_primitive} * state.geometry_stride +
indices_per_primitive)) *
sizeof(u32) <=
MAX_SCRATCH_SIZE};
const bool runs{total_primitives != 0 && fits};
if (!fits) {
LOG_WARNING(Render_Vulkan, "Draw of {} vertices and {} instances is too large to run as "
"compute", vertex_count, instance_count);
}
// Invocations past the end of the draw write into the padding of the buffers
const u32 vertex_invocations{
runs ? Common::AlignUp(static_cast<u32>(total_vertices), VTG_LOCAL_SIZE) : VTG_LOCAL_SIZE};
const u32 geometry_invocations{
runs ? Common::AlignUp(static_cast<u32>(total_primitives), VTG_LOCAL_SIZE)
: VTG_LOCAL_SIZE};
const u32 output_vertices{geometry_invocations * vertices_per_primitive};
const u32 output_indices{geometry_invocations * indices_per_primitive};
const bool emits_points{state.geometry_bindings.output_topology ==
Shader::OutputTopology::PointList};
const VkDeviceSize vertex_size{u64{vertex_invocations} * state.vertex_stride * sizeof(u32)};
// One more vertex is reserved, used by the indices of points that are not emitted
const VkDeviceSize geometry_size{u64{output_vertices + 1} * state.geometry_stride *
sizeof(u32)};
const VkDeviceSize index_size{u64{output_indices} * sizeof(u32)};
const auto [vertex_buffer, vertex_offset]{state.scratch->Allocate(vertex_size)};
const auto [geometry_buffer, geometry_offset]{state.scratch->Allocate(geometry_size)};
const auto [index_buffer, index_offset]{state.scratch->Allocate(index_size)};
// Describe the draw to the stages running as compute
const StagingBufferRef info_ref{state.staging_pool->Request(
DrawInfo::NUM_WORDS * sizeof(u32), MemoryUsage::Upload)};
const VtgState::Buffer info_buffer{info_ref.buffer, info_ref.offset,
DrawInfo::NUM_WORDS * sizeof(u32)};
std::array<u32, DrawInfo::NUM_WORDS> words{};
const VkDeviceSize alignment{
std::max<VkDeviceSize>(device.GetStorageBufferAlignment(), sizeof(u32))};
// Guest buffers are bound from an aligned offset, the remainder is added by the shader
const auto guest_buffer{[&](const std::tuple<Buffer*, u32, u32>& binding, u32& misalignment) {
const auto [buffer, offset, size]{binding};
misalignment = 0;
if (!buffer || buffer->Handle() == VK_NULL_HANDLE) {
return info_buffer;
}
const VkDeviceSize aligned{Common::AlignDown(VkDeviceSize{offset}, alignment)};
misalignment = static_cast<u32>(offset - aligned);
const VkDeviceSize wanted{Common::AlignUp(VkDeviceSize{size} + misalignment, 8)};
const VkDeviceSize available{
Common::AlignDown(VkDeviceSize{buffer->SizeBytes()} - aligned, sizeof(u32))};
return VtgState::Buffer{buffer->Handle(), aligned, std::min(wanted, available)};
}};
static_vector<VtgState::Buffer, 3 + DrawInfo::NUM_ATTRIBUTES> vertex_buffers;
std::array<u32, Maxwell::NumVertexArrays> misalignments{};
u32 index_misalignment{};
const u32 index_bytes{is_indexed ? draw_state.index_buffer.FormatSizeInBytes() : 0};
vertex_buffers.push_back(info_buffer);
vertex_buffers.push_back({vertex_buffer, vertex_offset, vertex_size});
vertex_buffers.push_back(is_indexed
? guest_buffer(buffer_cache.IndexBufferBinding(),
index_misalignment)
: info_buffer);
for (const auto& [stream, binding] : state.vertex_bindings.vertex_buffers) {
vertex_buffers.push_back(
guest_buffer(buffer_cache.VertexBufferBinding(stream), misalignments[stream]));
}
words[DrawInfo::VERTEX_COUNT] = vertex_count;
words[DrawInfo::INSTANCE_COUNT] = std::max(instance_count, 1U);
words[DrawInfo::BASE_VERTEX] =
is_indexed ? draw_state.base_index : draw_state.vertex_buffer.first;
words[DrawInfo::BASE_INSTANCE] = draw_state.base_instance;
words[DrawInfo::INDEX_SIZE] = index_bytes;
words[DrawInfo::INDEX_OFFSET] =
index_misalignment + (is_indexed ? draw_state.index_buffer.first * index_bytes : 0);
words[DrawInfo::PRIMITIVE_COUNT] = primitive_count;
words[DrawInfo::INDEX_MASK] =
index_bytes == 1 ? 0xffU : (index_bytes == 2 ? 0xffffU : 0xffffffffU);
for (size_t index = 0; index < DrawInfo::NUM_ATTRIBUTES; ++index) {
const auto& attribute{regs.vertex_attrib_format[index]};
const u32 stream{attribute.buffer};
u32* const attribute_words{&words[DrawInfo::ATTRIBUTES + index * DrawInfo::ATTRIBUTE_WORDS]};
attribute_words[DrawInfo::ATTRIBUTE_OFFSET] = misalignments[stream] + attribute.offset;
attribute_words[DrawInfo::ATTRIBUTE_STRIDE] = regs.vertex_streams[stream].stride;
attribute_words[DrawInfo::ATTRIBUTE_INSTANCED] =
regs.vertex_stream_instances.IsInstancingEnabled(stream) ? 1 : 0;
attribute_words[DrawInfo::ATTRIBUTE_DIVISOR] = regs.vertex_streams[stream].frequency;
}
std::memcpy(info_ref.mapped_span.data(), words.data(), sizeof(words));
const std::array<VtgState::Buffer, 4> geometry_buffers{{
info_buffer,
{vertex_buffer, vertex_offset, vertex_size},
{geometry_buffer, geometry_offset, geometry_size},
{index_buffer, index_offset, index_size},
}};
if (runs && !build_complete.load(std::memory_order_acquire)) {
scheduler.Record([this](vk::CommandBuffer) {
std::unique_lock lock{build_mutex};
build_condvar.wait(lock,
[this] { return build_complete.load(std::memory_order_acquire); });
});
}
const VideoCommon::SamplerId* samplers_it{samplers.data()};
const VideoCommon::ImageViewInOut* views_it{views.data()};
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();
buffer_cache.BindHostStageUniformAndStorageBuffers(guest_stage);
for (const VtgState::Buffer& buffer : buffers) {
guest_descriptor_queue.AddBuffer(buffer.buffer, buffer.offset, buffer.size);
}
buffer_cache.BindHostStageTextureBuffers(guest_stage);
RescalingPushConstant stage_rescaling;
PushImageDescriptors(texture_cache, guest_descriptor_queue, stage.guest_info,
stage_rescaling, samplers_it, views_it);
if (!runs) {
return;
}
const void* const descriptor_data{guest_descriptor_queue.UpdateData()};
const bool is_rescaling{!stage.guest_info.texture_descriptors.empty() ||
!stage.guest_info.image_descriptors.empty()};
scheduler.RequestOutsideRenderPassOperationContext();
scheduler.Record([this, &stage, descriptor_data, is_rescaling,
rescaling_data = stage_rescaling.Data(),
groups = invocations / VTG_LOCAL_SIZE](vk::CommandBuffer cmdbuf) {
if (!IsBuilt()) {
return;
}
cmdbuf.BindPipeline(VK_PIPELINE_BIND_POINT_COMPUTE, *stage.pipeline);
if (is_rescaling) {
cmdbuf.PushConstants(*stage.pipeline_layout, VK_SHADER_STAGE_COMPUTE_BIT,
RESCALING_LAYOUT_WORDS_OFFSET, sizeof(rescaling_data),
rescaling_data.data());
}
const VkDescriptorSet descriptor_set{stage.descriptor_allocator.Commit()};
device.GetLogical().UpdateDescriptorSet(descriptor_set,
*stage.descriptor_update_template,
descriptor_data);
cmdbuf.BindDescriptorSets(VK_PIPELINE_BIND_POINT_COMPUTE, *stage.pipeline_layout, 0,
descriptor_set, nullptr);
cmdbuf.Dispatch(groups, 1, 1);
static constexpr VkMemoryBarrier barrier{
.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER,
.pNext = nullptr,
.srcAccessMask = VK_ACCESS_SHADER_WRITE_BIT,
.dstAccessMask = VK_ACCESS_SHADER_READ_BIT | VK_ACCESS_INDEX_READ_BIT,
};
cmdbuf.PipelineBarrier(VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT,
VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT |
VK_PIPELINE_STAGE_VERTEX_INPUT_BIT |
VK_PIPELINE_STAGE_VERTEX_SHADER_BIT,
0, barrier);
});
}};
run_stage(VtgPipelineInfo::VERTEX, 0, vertex_buffers, vertex_invocations);
if (runs) {
// Index slots that are not written end the strip, or draw a discarded vertex as a point
const u32 unused_index{emits_points ? output_vertices : Shader::VTG_RESTART_INDEX};
const VkDeviceSize sentinel_offset{geometry_offset +
u64{output_vertices} * state.geometry_stride *
sizeof(u32)};
const VkDeviceSize sentinel_size{u64{state.geometry_stride} * sizeof(u32)};
scheduler.RequestOutsideRenderPassOperationContext();
scheduler.Record([index_buffer = index_buffer, index_offset = index_offset, index_size,
unused_index, geometry_buffer = geometry_buffer, sentinel_offset,
sentinel_size, emits_points](vk::CommandBuffer cmdbuf) {
static constexpr u32 QUIET_NAN = 0x7fc00000;
cmdbuf.FillBuffer(index_buffer, index_offset, index_size, unused_index);
if (emits_points) {
cmdbuf.FillBuffer(geometry_buffer, sentinel_offset, sentinel_size, QUIET_NAN);
}
static constexpr VkMemoryBarrier barrier{
.sType = VK_STRUCTURE_TYPE_MEMORY_BARRIER,
.pNext = nullptr,
.srcAccessMask = VK_ACCESS_TRANSFER_WRITE_BIT,
.dstAccessMask = VK_ACCESS_SHADER_WRITE_BIT,
};
cmdbuf.PipelineBarrier(VK_PIPELINE_STAGE_TRANSFER_BIT,
VK_PIPELINE_STAGE_COMPUTE_SHADER_BIT, 0, barrier);
});
}
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.AddBuffer(geometry_buffer, geometry_offset, geometry_size);
RescalingPushConstant rescaling;
RenderAreaPushConstant render_area;
buffer_cache.BindHostStageBuffers(FRAGMENT_STAGE);
PushImageDescriptors(texture_cache, guest_descriptor_queue, stage_infos[FRAGMENT_STAGE],
rescaling, samplers_it, views_it);
if (stage_infos[FRAGMENT_STAGE].uses_render_area) {
render_area.uses_render_area = true;
render_area.words = {static_cast<float>(regs.surface_clip.width),
static_cast<float>(regs.surface_clip.height)};
}
texture_cache.UpdateRenderTargets(false);
texture_cache.CheckFeedbackLoop(views);
ConfigureDraw(rescaling, render_area);
if (runs) {
state.index_buffer = index_buffer;
state.index_offset = index_offset;
state.index_count = static_cast<u32>(total_primitives) * indices_per_primitive;
}
}
void GraphicsPipeline::ConfigureDraw(const RescalingPushConstant& rescaling,
const RenderAreaPushConstant& render_area) {
const Framebuffer* const framebuffer = texture_cache.GetFramebuffer();
@ -732,12 +1209,28 @@ void GraphicsPipeline::MakePipeline(VkRenderPass render_pass) {
input_assembly_topology = VK_PRIMITIVE_TOPOLOGY_PATCH_LIST;
}
}
if (vtg) {
// The geometry stage running as compute writes strips cut by the restart index
switch (vtg->geometry_bindings.output_topology) {
case Shader::OutputTopology::PointList:
input_assembly_topology = VK_PRIMITIVE_TOPOLOGY_POINT_LIST;
break;
case Shader::OutputTopology::LineStrip:
input_assembly_topology = VK_PRIMITIVE_TOPOLOGY_LINE_STRIP;
break;
case Shader::OutputTopology::TriangleStrip:
input_assembly_topology = VK_PRIMITIVE_TOPOLOGY_TRIANGLE_STRIP;
break;
}
}
const bool vtg_restart{vtg && input_assembly_topology != VK_PRIMITIVE_TOPOLOGY_POINT_LIST};
const VkPipelineInputAssemblyStateCreateInfo input_assembly_ci{
.sType = VK_STRUCTURE_TYPE_PIPELINE_INPUT_ASSEMBLY_STATE_CREATE_INFO,
.pNext = nullptr,
.flags = 0,
.topology = input_assembly_topology,
.primitiveRestartEnable =
vtg_restart ||
dynamic.primitive_restart_enable != 0 &&
((input_assembly_topology != VK_PRIMITIVE_TOPOLOGY_PATCH_LIST &&
device.IsTopologyListPrimitiveRestartSupported()) ||

66
src/yuzu_video_core/renderer_vulkan/vk_graphics_pipeline.h
View File

@ -7,11 +7,15 @@
#include <array>
#include <atomic>
#include <condition_variable>
#include <memory>
#include <mutex>
#include <tuple>
#include <type_traits>
#include <utility>
#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/vk_buffer_cache.h"
@ -65,6 +69,47 @@ class RescalingPushConstant;
class RenderAreaPushConstant;
class Scheduler;
class StagingBufferPool;
/// Memory for the buffers shared between the stages of the draws running as compute
class VtgScratchAllocator {
public:
explicit VtgScratchAllocator(const Device& device, Scheduler& scheduler,
StagingBufferPool& staging_pool);
/// Returns a buffer and the offset of a region that lives until the current tick completes
[[nodiscard]] std::pair<VkBuffer, VkDeviceSize> Allocate(VkDeviceSize size);
private:
Scheduler& scheduler;
StagingBufferPool& staging_pool;
VkDeviceSize alignment;
VkBuffer buffer{};
VkDeviceSize base{};
VkDeviceSize capacity{};
VkDeviceSize used{};
u64 tick{};
};
/// Vertex and geometry stages of a pipeline running as compute before the draw
struct VtgPipelineInfo {
static constexpr size_t VERTEX = 0;
static constexpr size_t GEOMETRY = 1;
static constexpr size_t NUM_COMPUTE_STAGES = 2;
std::array<vk::ShaderModule, NUM_COMPUTE_STAGES> modules;
/// Shader information of each stage, including the storage buffers added to it
std::array<Shader::Info, NUM_COMPUTE_STAGES> infos;
Shader::VtgVertexBindings vertex_bindings;
Shader::VtgGeometryBindings geometry_bindings;
Shader::VtgTopology topology{};
/// Words of a vertex written by each stage
u32 vertex_stride{};
u32 geometry_stride{};
VtgScratchAllocator* scratch{};
StagingBufferPool* staging_pool{};
};
class GraphicsPipeline {
static constexpr size_t NUM_STAGES = Tegra::Engines::Maxwell3D::Regs::MaxShaderStage;
@ -76,7 +121,9 @@ public:
GuestDescriptorQueue& guest_descriptor_queue, Common::ThreadWorker* worker_thread,
PipelineStatistics* pipeline_statistics, RenderPassCache& render_pass_cache,
const GraphicsPipelineCacheKey& key, std::array<vk::ShaderModule, NUM_STAGES> stages,
const std::array<const Shader::Info*, NUM_STAGES>& infos, bool rasterizes_lines);
const std::array<const Shader::Info*, NUM_STAGES>& infos, bool rasterizes_lines,
std::unique_ptr<VtgPipelineInfo> vtg_info = nullptr);
~GraphicsPipeline();
GraphicsPipeline& operator=(GraphicsPipeline&&) noexcept = delete;
GraphicsPipeline(GraphicsPipeline&&) noexcept = delete;
@ -118,6 +165,18 @@ public:
return [](GraphicsPipeline* pl, bool is_indexed) { pl->ConfigureImpl<Spec>(is_indexed); };
}
/// Returns true when the vertex and geometry stages run as compute before the draw
[[nodiscard]] bool IsVtgAsCompute() const noexcept {
return vtg != nullptr;
}
/// Sets the number of instances of the next draw, zero when it cannot run as compute
void SetVtgInstanceCount(u32 instance_count) noexcept;
/// Index buffer, its offset and the number of indices drawing the output of the stages
/// running as compute. The count is zero when there is nothing to draw.
[[nodiscard]] std::tuple<VkBuffer, VkDeviceSize, u32> VtgDraw() const noexcept;
void SetEngine(Tegra::Engines::Maxwell3D* maxwell3d_, Tegra::MemoryManager* gpu_memory_) {
maxwell3d = maxwell3d_;
gpu_memory = gpu_memory_;
@ -127,6 +186,8 @@ private:
template <typename Spec>
void ConfigureImpl(bool is_indexed);
void ConfigureVtg(bool is_indexed);
void ConfigureDraw(const RescalingPushConstant& rescaling,
const RenderAreaPushConstant& render_are);
@ -171,6 +232,9 @@ private:
std::atomic_bool build_complete{false};
bool uses_push_descriptor{false};
const bool rasterizes_lines;
struct VtgState;
std::unique_ptr<VtgState> vtg;
};
} // namespace Vulkan

262
src/yuzu_video_core/renderer_vulkan/vk_pipeline_cache.cpp
View File

@ -17,6 +17,7 @@
#include "yuzu_shader_recompiler/environment.h"
#include "yuzu_shader_recompiler/frontend/maxwell/control_flow.h"
#include "yuzu_shader_recompiler/frontend/maxwell/translate_program.h"
#include "yuzu_shader_recompiler/ir_opt/passes.h"
#include "yuzu_shader_recompiler/program_header.h"
#include "yuzu_video_core/engines/kepler_compute.h"
#include "yuzu_video_core/engines/maxwell_3d.h"
@ -45,6 +46,7 @@ namespace {
using Shader::Backend::SPIRV::EmitSPIRV;
using Shader::Maxwell::ConvertLegacyToGeneric;
using Shader::Maxwell::GenerateGeometryPassthrough;
using Shader::Maxwell::GenerateVtgPassthroughVertex;
using Shader::Maxwell::MergeDualVertexPrograms;
using Shader::Maxwell::TranslateProgram;
using VideoCommon::ComputeEnvironment;
@ -293,6 +295,139 @@ size_t GetTotalPipelineWorkers() {
#endif
}
std::optional<Shader::VtgTopology> MaxwellToVtgTopology(Maxwell::PrimitiveTopology topology) {
switch (topology) {
case Maxwell::PrimitiveTopology::Points:
return Shader::VtgTopology::Points;
case Maxwell::PrimitiveTopology::Lines:
return Shader::VtgTopology::Lines;
case Maxwell::PrimitiveTopology::LineLoop:
return Shader::VtgTopology::LineLoop;
case Maxwell::PrimitiveTopology::LineStrip:
return Shader::VtgTopology::LineStrip;
case Maxwell::PrimitiveTopology::Triangles:
return Shader::VtgTopology::Triangles;
case Maxwell::PrimitiveTopology::TriangleStrip:
return Shader::VtgTopology::TriangleStrip;
case Maxwell::PrimitiveTopology::TriangleFan:
case Maxwell::PrimitiveTopology::Polygon:
return Shader::VtgTopology::TriangleFan;
case Maxwell::PrimitiveTopology::Quads:
return Shader::VtgTopology::Quads;
case Maxwell::PrimitiveTopology::QuadStrip:
return Shader::VtgTopology::QuadStrip;
case Maxwell::PrimitiveTopology::LinesAdjacency:
return Shader::VtgTopology::LinesAdjacency;
case Maxwell::PrimitiveTopology::LineStripAdjacency:
return Shader::VtgTopology::LineStripAdjacency;
case Maxwell::PrimitiveTopology::TrianglesAdjacency:
return Shader::VtgTopology::TrianglesAdjacency;
case Maxwell::PrimitiveTopology::TriangleStripAdjacency:
return Shader::VtgTopology::TriangleStripAdjacency;
case Maxwell::PrimitiveTopology::Patches:
break;
}
return std::nullopt;
}
/// Describes how a vertex shader running as compute reads its attributes from the vertex buffers
Shader::VtgVertexInputs MakeVtgVertexInputs(const FixedPipelineState& state) {
using Size = Maxwell::VertexAttribute::Size;
using Type = Maxwell::VertexAttribute::Type;
Shader::VtgVertexInputs inputs{};
for (size_t index = 0; index < inputs.size(); ++index) {
const FixedPipelineState::VertexAttribute& attribute{state.attributes[index]};
Shader::VtgVertexAttribute& input{inputs[index]};
if (attribute.enabled == 0) {
continue;
}
input.buffer = static_cast<u8>(attribute.buffer.Value());
const auto set_size{[&input](u8 components, u8 bits) {
input.components = components;
input.bits = bits;
}};
switch (attribute.Size()) {
case Size::Size_R32_G32_B32_A32:
set_size(4, 32);
break;
case Size::Size_R32_G32_B32:
set_size(3, 32);
break;
case Size::Size_R32_G32:
set_size(2, 32);
break;
case Size::Size_R32:
set_size(1, 32);
break;
case Size::Size_R16_G16_B16_A16:
set_size(4, 16);
break;
case Size::Size_R16_G16_B16:
set_size(3, 16);
break;
case Size::Size_R16_G16:
set_size(2, 16);
break;
case Size::Size_R16:
set_size(1, 16);
break;
case Size::Size_R8_G8_B8_A8:
case Size::Size_X8_B8_G8_R8:
set_size(4, 8);
break;
case Size::Size_R8_G8_B8:
set_size(3, 8);
break;
case Size::Size_R8_G8:
case Size::Size_G8_R8:
set_size(2, 8);
break;
case Size::Size_R8:
case Size::Size_A8:
set_size(1, 8);
break;
case Size::Size_A2_B10_G10_R10:
set_size(4, 0);
input.packed_offsets = {0, 10, 20, 30};
input.packed_bits = {10, 10, 10, 2};
break;
case Size::Size_B10_G11_R11:
set_size(3, 0);
input.packed_offsets = {0, 11, 22, 0};
input.packed_bits = {11, 11, 10, 0};
break;
default:
continue;
}
switch (attribute.Type()) {
case Type::SNorm:
input.type = Shader::VtgComponentType::SNorm;
break;
case Type::UNorm:
input.type = Shader::VtgComponentType::UNorm;
break;
case Type::SInt:
input.type = Shader::VtgComponentType::SInt;
break;
case Type::UInt:
input.type = Shader::VtgComponentType::UInt;
break;
case Type::UScaled:
input.type = Shader::VtgComponentType::UScaled;
break;
case Type::SScaled:
input.type = Shader::VtgComponentType::SScaled;
break;
case Type::Float:
input.type = Shader::VtgComponentType::Float;
break;
default:
break;
}
}
return inputs;
}
} // Anonymous namespace
size_t ComputePipelineCacheKey::Hash() const noexcept {
@ -318,8 +453,10 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
DescriptorPool& descriptor_pool_,
GuestDescriptorQueue& guest_descriptor_queue_,
RenderPassCache& render_pass_cache_, BufferCache& buffer_cache_,
TextureCache& texture_cache_, VideoCore::ShaderNotify& shader_notify_)
TextureCache& texture_cache_, VideoCore::ShaderNotify& shader_notify_,
StagingBufferPool& staging_pool_)
: VideoCommon::ShaderCache{device_memory_}, device{device_}, scheduler{scheduler_},
staging_pool{staging_pool_}, vtg_scratch{device_, scheduler_, staging_pool_},
descriptor_pool{descriptor_pool_}, guest_descriptor_queue{guest_descriptor_queue_},
render_pass_cache{render_pass_cache_}, buffer_cache{buffer_cache_},
texture_cache{texture_cache_}, shader_notify{shader_notify_},
@ -393,6 +530,10 @@ PipelineCache::PipelineCache(Tegra::MaxwellDeviceMemoryManager& device_memory_,
device.IsNvidia() && device.GetNvidiaArch() <= NvidiaArchitecture::Arch_Pascal,
.min_ssbo_alignment = device.GetStorageBufferAlignment(),
.max_user_clip_distances = device.GetMaxUserClipDistances(),
// Metal refuses shaders over its limit, the other drivers build them
.max_compute_shared_memory_size = device.GetDriverID() == VK_DRIVER_ID_MOLTENVK
? device.GetMaxComputeSharedMemorySize()
: 0,
};
host_info = Shader::HostTranslateInfo{
@ -722,11 +863,102 @@ std::unique_ptr<GraphicsPipeline> PipelineCache::CreateGraphicsPipeline(
if (skip_build && (*skip_build)()) {
return nullptr;
}
// Without host geometry shaders, a geometry shader that only selects the layer of the
// triangle it forwards is folded into the vertex shader
std::optional<Shader::IR::Attribute> forwarded_layer;
const size_t geometry_index{static_cast<size_t>(Maxwell::ShaderType::Geometry)};
if (!device.IsGeometryShaderSupported() && device.IsExtShaderViewportIndexLayerSupported() &&
key.unique_hashes[geometry_index] != 0 && key.unique_hashes[2] == 0 &&
key.unique_hashes[3] == 0 && uses_vertex_b) {
const u32 input_vertices{Shader::InputTopologyVertices::vertices(
MakeRuntimeInfo(programs, key, programs[geometry_index], nullptr).input_topology)};
forwarded_layer = Shader::Optimization::FindForwardedLayerAttribute(
programs[geometry_index], input_vertices);
if (forwarded_layer) {
programs[1].info.stores.Set(Shader::IR::Attribute::Layer);
}
}
std::array<const Shader::Info*, Maxwell::MaxShaderStage> infos{};
std::array<vk::ShaderModule, Maxwell::MaxShaderStage> modules;
const Shader::IR::Program* previous_stage{};
Shader::Backend::Bindings binding;
// Without host geometry shaders, the vertex and geometry shaders run as compute shaders
// before the draw and a generated vertex shader feeds their output to the rasterizer
std::unique_ptr<VtgPipelineInfo> vtg_info;
Shader::IR::Program vtg_passthrough;
const std::optional<Shader::VtgTopology> vtg_topology{
MaxwellToVtgTopology(key.state.topology)};
if (!device.IsGeometryShaderSupported() && !forwarded_layer && vtg_topology &&
key.unique_hashes[geometry_index] != 0 && key.unique_hashes[2] == 0 &&
key.unique_hashes[3] == 0 && uses_vertex_b && key.state.dynamic_vertex_input == 0 &&
key.state.xfb_enabled == 0) {
Shader::IR::Program& vertex{programs[1]};
Shader::IR::Program& geometry{programs[geometry_index]};
const auto vertex_runtime{MakeRuntimeInfo(programs, key, vertex, nullptr)};
ConvertLegacyToGeneric(vertex, vertex_runtime);
const auto geometry_runtime{MakeRuntimeInfo(programs, key, geometry, &vertex)};
ConvertLegacyToGeneric(geometry, geometry_runtime);
Shader::VaryingState geometry_stores{geometry.info.stores};
if (!device.IsExtShaderViewportIndexLayerSupported()) {
geometry_stores.Set(Shader::IR::Attribute::Layer, false);
geometry_stores.Set(Shader::IR::Attribute::ViewportIndex, false);
}
const Shader::VtgVaryingLayout vertex_layout{vertex.info.stores};
const Shader::VtgVaryingLayout geometry_layout{geometry_stores};
const u32 input_vertices{
Shader::InputTopologyVertices::vertices(geometry_runtime.input_topology)};
const Shader::OutputTopology geometry_topology{geometry.output_topology};
auto info{std::make_unique<VtgPipelineInfo>()};
if (Shader::Optimization::VtgVertexToCompute(vertex, MakeVtgVertexInputs(key.state),
vertex_layout, info->vertex_bindings)) {
if (!Shader::Optimization::VtgGeometryToCompute(geometry, *vtg_topology,
input_vertices, vertex_layout,
geometry_layout,
info->geometry_bindings)) {
LOG_ERROR(Render_Vulkan, "Geometry shader {:016x} cannot run as compute",
key.unique_hashes[geometry_index]);
return nullptr;
}
const std::array<Shader::IR::Program*, VtgPipelineInfo::NUM_COMPUTE_STAGES> stages{
&vertex, &geometry};
const std::array<const Shader::RuntimeInfo*, VtgPipelineInfo::NUM_COMPUTE_STAGES>
runtimes{&vertex_runtime, &geometry_runtime};
for (size_t stage = 0; stage < stages.size(); ++stage) {
Shader::Backend::Bindings compute_binding;
const std::vector<u32> code{
EmitSPIRV(profile, *runtimes[stage], *stages[stage], compute_binding)};
device.SaveShader(code);
info->modules[stage] = BuildShader(device, code);
info->infos[stage] = stages[stage]->info;
}
info->topology = *vtg_topology;
info->vertex_stride = vertex_layout.stride;
info->geometry_stride = geometry_layout.stride;
info->scratch = &vtg_scratch;
info->staging_pool = &staging_pool;
vtg_passthrough = GenerateVtgPassthroughVertex(pools.inst, pools.block, geometry,
geometry_stores, geometry_layout);
auto runtime_info{MakeRuntimeInfo(programs, key, vtg_passthrough, nullptr)};
runtime_info.convert_depth_mode = key.state.ndc_minus_one_to_one != 0;
if (geometry_topology == Shader::OutputTopology::PointList) {
runtime_info.fixed_state_point_size = Common::BitCast<float>(key.state.point_size);
}
const std::vector<u32> code{EmitSPIRV(profile, runtime_info, vtg_passthrough, binding)};
device.SaveShader(code);
modules[0] = BuildShader(device, code);
infos[0] = &vtg_passthrough.info;
previous_stage = &vtg_passthrough;
vtg_info = std::move(info);
LOG_INFO(Render_Vulkan, "Running vertex shader {:016x} and geometry shader {:016x} "
"as compute",
key.unique_hashes[1], key.unique_hashes[geometry_index]);
}
}
for (size_t index = uses_vertex_a && uses_vertex_b ? 1 : 0; index < Maxwell::MaxShaderProgram;
++index) {
const bool is_emulated_stage = layer_source_program != nullptr &&
@ -734,13 +966,23 @@ std::unique_ptr<GraphicsPipeline> PipelineCache::CreateGraphicsPipeline(
if (key.unique_hashes[index] == 0 && !is_emulated_stage) {
continue;
}
if (forwarded_layer && index == geometry_index) {
continue;
}
if (vtg_info && (index == 1 || index == geometry_index)) {
continue;
}
UNIMPLEMENTED_IF(index == 0);
Shader::IR::Program& program{programs[index]};
const size_t stage_index{index - 1};
infos[stage_index] = &program.info;
const auto runtime_info{MakeRuntimeInfo(programs, key, program, previous_stage)};
auto runtime_info{MakeRuntimeInfo(programs, key, program, previous_stage)};
if (forwarded_layer && program.stage == Shader::Stage::VertexB) {
runtime_info.forwarded_layer_attribute = forwarded_layer;
runtime_info.convert_depth_mode = key.state.ndc_minus_one_to_one != 0;
}
ConvertLegacyToGeneric(program, runtime_info);
const std::vector<u32> code{EmitSPIRV(profile, runtime_info, program, binding)};
device.SaveShader(code);
@ -757,14 +999,15 @@ std::unique_ptr<GraphicsPipeline> PipelineCache::CreateGraphicsPipeline(
Maxwell::Tessellation::DomainType::Isolines};
const auto output_topology{GetRasterizationOutputTopology(
key.state.topology, has_tessellation, tessellation_isolines,
modules[3] ? std::optional{programs[4].output_topology} : std::nullopt)};
vtg_info ? std::optional{vtg_info->geometry_bindings.output_topology}
: (modules[3] ? std::optional{programs[4].output_topology} : std::nullopt))};
const bool rasterizes_lines{RasterizesLines(
output_topology, FixedPipelineState::UnpackPolygonMode(key.state.polygon_mode))};
Common::ThreadWorker* const thread_worker{build_in_parallel ? &workers : nullptr};
return std::make_unique<GraphicsPipeline>(
scheduler, buffer_cache, texture_cache, vulkan_pipeline_cache, &shader_notify, device,
descriptor_pool, guest_descriptor_queue, thread_worker, statistics, render_pass_cache, key,
std::move(modules), infos, rasterizes_lines);
std::move(modules), infos, rasterizes_lines, std::move(vtg_info));
} catch (const Shader::Exception& exception) {
auto hash = key.Hash();
@ -878,6 +1121,17 @@ std::unique_ptr<ComputePipeline> PipelineCache::CreateComputePipeline(
}
auto program{TranslateProgram(pools.inst, pools.block, env, cfg, host_info)};
// Only where the host enforces its limit
const u32 max_shared_memory{profile.max_compute_shared_memory_size};
if (max_shared_memory != 0 && program.shared_memory_size > max_shared_memory &&
!profile.support_explicit_workgroup_layout) {
const u32 requested_shared_memory{program.shared_memory_size};
if (Shader::Optimization::TryPackSharedMemory16To12(program, max_shared_memory)) {
LOG_INFO(Render_Vulkan,
"Compute shader {:#016x}: compacted shared memory from {} to {} bytes",
key.unique_hash, requested_shared_memory, program.shared_memory_size);
}
}
const std::vector<u32> code{EmitSPIRV(profile, program)};
device.SaveShader(code);
vk::ShaderModule spv_module{BuildShader(device, code)};

6
src/yuzu_video_core/renderer_vulkan/vk_pipeline_cache.h
View File

@ -84,6 +84,7 @@ class Device;
class PipelineStatistics;
class RenderPassCache;
class Scheduler;
class StagingBufferPool;
using VideoCommon::ShaderInfo;
@ -105,7 +106,8 @@ public:
Scheduler& scheduler, DescriptorPool& descriptor_pool,
GuestDescriptorQueue& guest_descriptor_queue,
RenderPassCache& render_pass_cache, BufferCache& buffer_cache,
TextureCache& texture_cache, VideoCore::ShaderNotify& shader_notify_);
TextureCache& texture_cache, VideoCore::ShaderNotify& shader_notify_,
StagingBufferPool& staging_pool);
~PipelineCache();
[[nodiscard]] GraphicsPipeline* CurrentGraphicsPipeline();
@ -144,6 +146,8 @@ private:
const Device& device;
Scheduler& scheduler;
StagingBufferPool& staging_pool;
VtgScratchAllocator vtg_scratch;
DescriptorPool& descriptor_pool;
GuestDescriptorQueue& guest_descriptor_queue;
RenderPassCache& render_pass_cache;

50
src/yuzu_video_core/renderer_vulkan/vk_rasterizer.cpp
View File

@ -189,7 +189,8 @@ RasterizerVulkan::RasterizerVulkan(Core::Frontend::EmuWindow& emu_window_, Tegra
staging_pool, compute_pass_descriptor_queue, descriptor_pool),
query_cache(gpu, *this, device_memory, query_cache_runtime),
pipeline_cache(device_memory, device, scheduler, descriptor_pool, guest_descriptor_queue,
render_pass_cache, buffer_cache, texture_cache, gpu.ShaderNotify()),
render_pass_cache, buffer_cache, texture_cache, gpu.ShaderNotify(),
staging_pool),
accelerate_dma(buffer_cache, texture_cache, scheduler),
fence_manager(*this, gpu, texture_cache, buffer_cache, query_cache, device, scheduler),
wfi_event(device.GetLogical().CreateEvent()) {
@ -215,6 +216,7 @@ void RasterizerVulkan::PrepareDraw(bool is_indexed, Func&& draw_func) {
std::scoped_lock lock{buffer_cache.mutex, texture_cache.mutex};
// update engine as channel may be different.
pipeline->SetEngine(maxwell3d, gpu_memory);
pipeline->SetVtgInstanceCount(std::exchange(vtg_instance_count, 0));
pipeline->Configure(is_indexed);
UpdateDynamicStates();
@ -225,8 +227,38 @@ void RasterizerVulkan::PrepareDraw(bool is_indexed, Func&& draw_func) {
draw_func(pipeline);
}
void RasterizerVulkan::DrawVtgAsCompute(const GraphicsPipeline* pipeline) {
// The stages running as compute wrote the strips to draw into an index buffer
const auto [index_buffer, index_offset, index_count]{pipeline->VtgDraw()};
if (index_count == 0) {
return;
}
const bool dynamic_restart{device.IsExtExtendedDynamicState2Supported()};
if (dynamic_restart) {
// The guest state is restored by the next draw
maxwell3d->dirty.flags[Dirty::PrimitiveRestartEnable] = true;
maxwell3d->dirty.flags[Dirty::StateEnable] = true;
}
scheduler.Record([pipeline, index_buffer = index_buffer, index_offset = index_offset,
index_count = index_count, dynamic_restart](vk::CommandBuffer cmdbuf) {
if (!pipeline->IsBuilt()) {
return;
}
if (dynamic_restart) {
cmdbuf.SetPrimitiveRestartEnableEXT(true);
}
cmdbuf.BindIndexBuffer(index_buffer, index_offset, VK_INDEX_TYPE_UINT32);
cmdbuf.DrawIndexed(index_count, 1, 0, 0, 0);
});
}
void RasterizerVulkan::Draw(bool is_indexed, u32 instance_count) {
vtg_instance_count = instance_count;
PrepareDraw(is_indexed, [this, is_indexed, instance_count](const GraphicsPipeline* pipeline) {
if (pipeline->IsVtgAsCompute()) {
DrawVtgAsCompute(pipeline);
return;
}
const auto& draw_state = maxwell3d->draw_manager->GetDrawState();
const u32 num_instances{instance_count};
const DrawParams draw_params{MakeDrawParams(draw_state, num_instances, is_indexed)};
@ -250,6 +282,10 @@ void RasterizerVulkan::DrawIndirect() {
const auto& params = maxwell3d->draw_manager->GetIndirectParams();
buffer_cache.SetDrawIndirect(&params);
PrepareDraw(params.is_indexed, [this, &params](const GraphicsPipeline* pipeline) {
if (pipeline->IsVtgAsCompute()) {
// The stages running as compute need the draw parameters on the CPU
return;
}
const auto indirect_buffer = buffer_cache.GetDrawIndirectBuffer();
const auto& buffer = indirect_buffer.first;
const auto& offset = indirect_buffer.second;
@ -539,15 +575,23 @@ void RasterizerVulkan::DispatchCompute() {
const auto [buffer, offset] =
buffer_cache.ObtainBuffer(*indirect_address, 12, sync_info, post_op);
scheduler.RequestOutsideRenderPassOperationContext();
scheduler.Record([indirect_buffer = buffer->Handle(),
scheduler.Record([pipeline, indirect_buffer = buffer->Handle(),
indirect_offset = offset](vk::CommandBuffer cmdbuf) {
if (pipeline->HasBuildFailed()) {
return;
}
cmdbuf.DispatchIndirect(indirect_buffer, indirect_offset);
});
return;
}
const std::array<u32, 3> dim{qmd.grid_dim_x, qmd.grid_dim_y, qmd.grid_dim_z};
scheduler.RequestOutsideRenderPassOperationContext();
scheduler.Record([dim](vk::CommandBuffer cmdbuf) { cmdbuf.Dispatch(dim[0], dim[1], dim[2]); });
scheduler.Record([pipeline, dim](vk::CommandBuffer cmdbuf) {
if (pipeline->HasBuildFailed()) {
return;
}
cmdbuf.Dispatch(dim[0], dim[1], dim[2]);
});
}
void RasterizerVulkan::ResetCounter(VideoCommon::QueryType type) {

4
src/yuzu_video_core/renderer_vulkan/vk_rasterizer.h
View File

@ -150,6 +150,8 @@ private:
template <typename Func>
void PrepareDraw(bool is_indexed, Func&&);
void DrawVtgAsCompute(const GraphicsPipeline* pipeline);
void FlushWork();
void UpdateDynamicStates();
@ -191,6 +193,8 @@ private:
Scheduler& scheduler;
StagingBufferPool staging_pool;
/// Instances of the draw being prepared, zero when it cannot run its stages as compute
u32 vtg_instance_count{};
DescriptorPool descriptor_pool;
GuestDescriptorQueue guest_descriptor_queue;
ComputePassDescriptorQueue compute_pass_descriptor_queue;

5
src/yuzu_video_core/vulkan_common/vulkan_device.h
View File

@ -520,6 +520,11 @@ public:
return extensions.depth_bias_control;
}
/// Returns true if the device supports geometry shaders.
bool IsGeometryShaderSupported() const {
return features.features.geometryShader;
}
/// Returns true if the device supports VK_EXT_shader_viewport_index_layer.
bool IsExtShaderViewportIndexLayerSupported() const {
return extensions.shader_viewport_index_layer;