mirrors/shadPS4
1
0
Fork 0
mirror of https://github.com/shadps4-emu/shadPS4.git synced 2025-08-04 16:32:39 +00:00
Code Issues Packages Projects Releases Wiki Activity

clang format

Browse Source
This commit is contained in:
IndecisiveTurtle 2025-02-14 12:16:34 +02:00
parent b05e6577d1
commit 67998c4a02
9 changed files with 75 additions and 47 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

12
src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp
View File

@ -165,9 +165,11 @@ using BufferAlias = EmitContext::BufferAlias;
Id EmitReadConst(EmitContext& ctx, IR::Inst* inst) {
const u32 flatbuf_off_dw = inst->Flags<u32>();
const auto& srt_flatbuf = ctx.buffers.back();
ASSERT(srt_flatbuf.binding >= 0 && flatbuf_off_dw > 0 && srt_flatbuf.buffer_type == BufferType::ReadConstUbo);
ASSERT(srt_flatbuf.binding >= 0 && flatbuf_off_dw > 0 &&
srt_flatbuf.buffer_type == BufferType::ReadConstUbo);
const auto [id, pointer_type] = srt_flatbuf[BufferAlias::U32];
const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, ctx.ConstU32(flatbuf_off_dw))};
const Id ptr{
ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, ctx.ConstU32(flatbuf_off_dw))};
return ctx.OpLoad(ctx.U32[1], ptr);
}
@ -416,8 +418,7 @@ static Id EmitLoadBufferB32xN(EmitContext& ctx, u32 handle, Id address) {
boost::container::static_vector<Id, N> ids;
for (u32 i = 0; i < N; i++) {
const Id index_i = ctx.OpIAdd(ctx.U32[1], index, ctx.ConstU32(i));
const Id ptr{
ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index_i)};
const Id ptr{ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index_i)};
ids.push_back(ctx.OpLoad(data_types[1], ptr));
}
return ctx.OpCompositeConstruct(data_types[N], ids);
@ -496,8 +497,7 @@ static void EmitStoreBufferB32xN(EmitContext& ctx, u32 handle, Id address, Id va
} else {
for (u32 i = 0; i < N; i++) {
const Id index_i = ctx.OpIAdd(ctx.U32[1], index, ctx.ConstU32(i));
const Id ptr =
ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index_i);
const Id ptr = ctx.OpAccessChain(pointer_type, id, ctx.u32_zero_value, index_i);
ctx.OpStore(ptr, ctx.OpCompositeExtract(data_types[1], value, i));
}
}

35
src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
View File

@ -248,8 +248,9 @@ void EmitContext::DefineWorkgroupIndex() {
const Id num_workgroups{OpLoad(U32[3], num_workgroups_id)};
const Id num_workgroups_x{OpCompositeExtract(U32[1], num_workgroups, 0)};
const Id num_workgroups_y{OpCompositeExtract(U32[1], num_workgroups, 1)};
workgroup_index_id = OpIAdd(U32[1], OpIAdd(U32[1], workgroup_x, OpIMul(U32[1], workgroup_y, num_workgroups_x)),
OpIMul(U32[1], workgroup_z, OpIMul(U32[1], num_workgroups_x, num_workgroups_y)));
workgroup_index_id =
OpIAdd(U32[1], OpIAdd(U32[1], workgroup_x, OpIMul(U32[1], workgroup_y, num_workgroups_x)),
OpIMul(U32[1], workgroup_z, OpIMul(U32[1], num_workgroups_x, num_workgroups_y)));
Name(workgroup_index_id, "workgroup_index");
}
@ -328,7 +329,8 @@ void EmitContext::DefineInputs() {
frag_depth = DefineVariable(F32[1], spv::BuiltIn::FragDepth, spv::StorageClass::Output);
}
if (info.loads.Get(IR::Attribute::IsFrontFace)) {
front_facing = DefineVariable(U1[1], spv::BuiltIn::FrontFacing, spv::StorageClass::Input);
front_facing =
DefineVariable(U1[1], spv::BuiltIn::FrontFacing, spv::StorageClass::Input);
}
if (profile.needs_manual_interpolation) {
gl_bary_coord_id =
@ -364,11 +366,14 @@ void EmitContext::DefineInputs() {
}
break;
case LogicalStage::Compute:
if (info.loads.GetAny(IR::Attribute::WorkgroupIndex) || info.loads.GetAny(IR::Attribute::WorkgroupId)) {
workgroup_id = DefineVariable(U32[3], spv::BuiltIn::WorkgroupId, spv::StorageClass::Input);
if (info.loads.GetAny(IR::Attribute::WorkgroupIndex) ||
info.loads.GetAny(IR::Attribute::WorkgroupId)) {
workgroup_id =
DefineVariable(U32[3], spv::BuiltIn::WorkgroupId, spv::StorageClass::Input);
}
if (info.loads.GetAny(IR::Attribute::WorkgroupIndex)) {
num_workgroups_id = DefineVariable(U32[3], spv::BuiltIn::NumWorkgroups, spv::StorageClass::Input);
num_workgroups_id =
DefineVariable(U32[3], spv::BuiltIn::NumWorkgroups, spv::StorageClass::Input);
}
if (info.loads.GetAny(IR::Attribute::LocalInvocationId)) {
local_invocation_id =
@ -625,7 +630,8 @@ EmitContext::BufferSpv EmitContext::DefineBuffer(bool is_storage, bool is_writte
: TypeArray(data_type, max_num_items)};
// Define block struct type. Don't perform decorations twice on the same Id.
const Id struct_type{TypeStruct(record_array_type)};
if (std::ranges::find(buf_type_ids, record_array_type.value, &Id::value) == buf_type_ids.end()) {
if (std::ranges::find(buf_type_ids, record_array_type.value, &Id::value) ==
buf_type_ids.end()) {
Decorate(record_array_type, spv::Decoration::ArrayStride, 1 << elem_shift);
Decorate(struct_type, spv::Decoration::Block);
MemberName(struct_type, 0, "data");
@ -668,16 +674,20 @@ void EmitContext::DefineBuffers() {
// Define aliases depending on the shader usage.
auto& spv_buffer = buffers.emplace_back(binding.buffer++, desc.buffer_type);
if (True(desc.used_types & IR::Type::U32)) {
spv_buffer[BufferAlias::U32] = DefineBuffer(is_storage, desc.is_written, 2, desc.buffer_type, U32[1]);
spv_buffer[BufferAlias::U32] =
DefineBuffer(is_storage, desc.is_written, 2, desc.buffer_type, U32[1]);
}
if (True(desc.used_types & IR::Type::F32)) {
spv_buffer[BufferAlias::F32] = DefineBuffer(is_storage, desc.is_written, 2, desc.buffer_type, F32[1]);
spv_buffer[BufferAlias::F32] =
DefineBuffer(is_storage, desc.is_written, 2, desc.buffer_type, F32[1]);
}
if (True(desc.used_types & IR::Type::U16)) {
spv_buffer[BufferAlias::U16] = DefineBuffer(is_storage, desc.is_written, 1, desc.buffer_type, U16);
spv_buffer[BufferAlias::U16] =
DefineBuffer(is_storage, desc.is_written, 1, desc.buffer_type, U16);
}
if (True(desc.used_types & IR::Type::U8)) {
spv_buffer[BufferAlias::U8] = DefineBuffer(is_storage, desc.is_written, 0, desc.buffer_type, U8);
spv_buffer[BufferAlias::U8] =
DefineBuffer(is_storage, desc.is_written, 0, desc.buffer_type, U8);
}
++binding.unified;
}
@ -835,7 +845,8 @@ void EmitContext::DefineSharedMemory() {
ASSERT(info.stage == Stage::Compute);
const u32 shared_memory_size = runtime_info.cs_info.shared_memory_size;
const u32 num_elements{Common::DivCeil(shared_memory_size, 4U)};
LOG_ERROR(Render_Recompiler, "Defined {:#x} num_elements = {}, shared_memory_size = {}", info.pgm_hash, num_elements, shared_memory_size);
LOG_ERROR(Render_Recompiler, "Defined {:#x} num_elements = {}, shared_memory_size = {}",
info.pgm_hash, num_elements, shared_memory_size);
const Id type{TypeArray(U32[1], ConstU32(num_elements))};
shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type);
shared_u32 = TypePointer(spv::StorageClass::Workgroup, U32[1]);

4
src/shader_recompiler/backend/spirv/spirv_emit_context.h
View File

@ -298,8 +298,8 @@ private:
SpirvAttribute GetAttributeInfo(AmdGpu::NumberFormat fmt, Id id, u32 num_components,
bool output);
BufferSpv DefineBuffer(bool is_storage, bool is_written, u32 elem_shift,
BufferType buffer_type, Id data_type);
BufferSpv DefineBuffer(bool is_storage, bool is_written, u32 elem_shift, BufferType buffer_type,
Id data_type);
Id DefineFloat32ToUfloatM5(u32 mantissa_bits, std::string_view name);
Id DefineUfloatM5ToFloat32(u32 mantissa_bits, std::string_view name);

6
src/shader_recompiler/frontend/translate/data_share.cpp
View File

@ -178,7 +178,8 @@ void Translator::DS_WRITE(int bit_size, bool is_signed, bool is_pair, bool strid
const IR::VectorReg data1{inst.src[2].code};
const u32 offset = (inst.control.ds.offset1 << 8u) + inst.control.ds.offset0;
if (info.stage == Stage::Fragment) {
ASSERT_MSG(!is_pair && bit_size == 32 && offset % 256 == 0, "Unexpected shared memory offset alignment: {}", offset);
ASSERT_MSG(!is_pair && bit_size == 32 && offset % 256 == 0,
"Unexpected shared memory offset alignment: {}", offset);
ir.SetVectorReg(GetScratchVgpr(offset), ir.GetVectorReg(data0));
return;
}
@ -229,7 +230,8 @@ void Translator::DS_READ(int bit_size, bool is_signed, bool is_pair, bool stride
IR::VectorReg dst_reg{inst.dst[0].code};
const u32 offset = (inst.control.ds.offset1 << 8u) + inst.control.ds.offset0;
if (info.stage == Stage::Fragment) {
ASSERT_MSG(!is_pair && bit_size == 32 && offset % 256 == 0, "Unexpected shared memory offset alignment: {}", offset);
ASSERT_MSG(!is_pair && bit_size == 32 && offset % 256 == 0,
"Unexpected shared memory offset alignment: {}", offset);
ir.SetVectorReg(dst_reg, ir.GetVectorReg(GetScratchVgpr(offset)));
return;
}

3
src/shader_recompiler/ir/passes/resource_tracking_pass.cpp
View File

@ -134,7 +134,8 @@ public:
u32 Add(const BufferResource& desc) {
const u32 index{Add(buffer_resources, desc, [&desc](const auto& existing) {
return desc.sharp_idx == existing.sharp_idx && desc.inline_cbuf == existing.inline_cbuf &&
return desc.sharp_idx == existing.sharp_idx &&
desc.inline_cbuf == existing.inline_cbuf &&
desc.buffer_type == existing.buffer_type;
})};
auto& buffer = buffer_resources[index];

15
src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp
View File

@ -67,15 +67,19 @@ static void EmitBarrierInMergeBlock(const IR::AbstractSyntaxNode::Data& data) {
static constexpr u32 GcnSubgroupSize = 64;
void SharedMemoryBarrierPass(IR::Program& program, const RuntimeInfo& runtime_info, const Profile& profile) {
void SharedMemoryBarrierPass(IR::Program& program, const RuntimeInfo& runtime_info,
const Profile& profile) {
if (program.info.stage != Stage::Compute) {
return;
}
const auto& cs_info = runtime_info.cs_info;
const u32 shared_memory_size = cs_info.shared_memory_size;
const u32 threadgroup_size = cs_info.workgroup_size[0] * cs_info.workgroup_size[1] * cs_info.workgroup_size[2];
// The compiler can only omit barriers when the local workgroup size is the same as the HW subgroup.
if (shared_memory_size == 0 || threadgroup_size != GcnSubgroupSize || !profile.needs_lds_barriers) {
const u32 threadgroup_size =
cs_info.workgroup_size[0] * cs_info.workgroup_size[1] * cs_info.workgroup_size[2];
// The compiler can only omit barriers when the local workgroup size is the same as the HW
// subgroup.
if (shared_memory_size == 0 || threadgroup_size != GcnSubgroupSize ||
!profile.needs_lds_barriers) {
return;
}
using Type = IR::AbstractSyntaxNode::Type;
@ -85,7 +89,8 @@ void SharedMemoryBarrierPass(IR::Program& program, const RuntimeInfo& runtime_in
--branch_depth;
continue;
}
// Check if branch depth is zero, we don't want to insert barrier in potentially divergent code.
// Check if branch depth is zero, we don't want to insert barrier in potentially divergent
// code.
if (node.type == Type::If && branch_depth++ == 0) {
EmitBarrierInMergeBlock(node.data);
continue;

24
src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp
View File

@ -28,7 +28,8 @@ static bool IsSharedAccess(const IR::Inst& inst) {
}
}
void SharedMemoryToStoragePass(IR::Program& program, const RuntimeInfo& runtime_info, const Profile& profile) {
void SharedMemoryToStoragePass(IR::Program& program, const RuntimeInfo& runtime_info,
const Profile& profile) {
if (program.info.stage != Stage::Compute) {
return;
}
@ -55,34 +56,41 @@ void SharedMemoryToStoragePass(IR::Program& program, const RuntimeInfo& runtime_
// Replace shared atomics first
switch (inst.GetOpcode()) {
case IR::Opcode::SharedAtomicAnd32:
inst.ReplaceUsesWithAndRemove(ir.BufferAtomicAnd(handle, inst.Arg(0), inst.Arg(1), {}));
inst.ReplaceUsesWithAndRemove(
ir.BufferAtomicAnd(handle, inst.Arg(0), inst.Arg(1), {}));
continue;
case IR::Opcode::SharedAtomicIAdd32:
inst.ReplaceUsesWithAndRemove(ir.BufferAtomicIAdd(handle, inst.Arg(0), inst.Arg(1), {}));
inst.ReplaceUsesWithAndRemove(
ir.BufferAtomicIAdd(handle, inst.Arg(0), inst.Arg(1), {}));
continue;
case IR::Opcode::SharedAtomicOr32:
inst.ReplaceUsesWithAndRemove(ir.BufferAtomicOr(handle, inst.Arg(0), inst.Arg(1), {}));
inst.ReplaceUsesWithAndRemove(
ir.BufferAtomicOr(handle, inst.Arg(0), inst.Arg(1), {}));
continue;
case IR::Opcode::SharedAtomicSMax32:
case IR::Opcode::SharedAtomicUMax32: {
const bool is_signed = inst.GetOpcode() == IR::Opcode::SharedAtomicSMax32;
inst.ReplaceUsesWithAndRemove(ir.BufferAtomicIMax(handle, inst.Arg(0), inst.Arg(1), is_signed, {}));
inst.ReplaceUsesWithAndRemove(
ir.BufferAtomicIMax(handle, inst.Arg(0), inst.Arg(1), is_signed, {}));
continue;
}
case IR::Opcode::SharedAtomicSMin32:
case IR::Opcode::SharedAtomicUMin32: {
const bool is_signed = inst.GetOpcode() == IR::Opcode::SharedAtomicSMin32;
inst.ReplaceUsesWithAndRemove(ir.BufferAtomicIMin(handle, inst.Arg(0), inst.Arg(1), is_signed, {}));
inst.ReplaceUsesWithAndRemove(
ir.BufferAtomicIMin(handle, inst.Arg(0), inst.Arg(1), is_signed, {}));
continue;
}
case IR::Opcode::SharedAtomicXor32:
inst.ReplaceUsesWithAndRemove(ir.BufferAtomicXor(handle, inst.Arg(0), inst.Arg(1), {}));
inst.ReplaceUsesWithAndRemove(
ir.BufferAtomicXor(handle, inst.Arg(0), inst.Arg(1), {}));
continue;
default:
break;
}
// Replace shared operations.
const IR::U32 offset = ir.IMul(ir.GetAttributeU32(IR::Attribute::WorkgroupIndex), ir.Imm32(shared_memory_size));
const IR::U32 offset = ir.IMul(ir.GetAttributeU32(IR::Attribute::WorkgroupIndex),
ir.Imm32(shared_memory_size));
const IR::U32 address = ir.IAdd(IR::U32{inst.Arg(0)}, offset);
switch (inst.GetOpcode()) {
case IR::Opcode::LoadSharedU32:

14
src/video_core/renderer_vulkan/vk_instance.cpp
View File

@ -203,14 +203,12 @@ std::string Instance::GetDriverVersionName() {
}
bool Instance::CreateDevice() {
const vk::StructureChain feature_chain =
physical_device
.getFeatures2<vk::PhysicalDeviceFeatures2, vk::PhysicalDeviceVulkan11Features,
vk::PhysicalDeviceVulkan12Features,
vk::PhysicalDeviceRobustness2FeaturesEXT,
vk::PhysicalDeviceExtendedDynamicState3FeaturesEXT,
vk::PhysicalDevicePrimitiveTopologyListRestartFeaturesEXT,
vk::PhysicalDevicePortabilitySubsetFeaturesKHR>();
const vk::StructureChain feature_chain = physical_device.getFeatures2<
vk::PhysicalDeviceFeatures2, vk::PhysicalDeviceVulkan11Features,
vk::PhysicalDeviceVulkan12Features, vk::PhysicalDeviceRobustness2FeaturesEXT,
vk::PhysicalDeviceExtendedDynamicState3FeaturesEXT,
vk::PhysicalDevicePrimitiveTopologyListRestartFeaturesEXT,
vk::PhysicalDevicePortabilitySubsetFeaturesKHR>();
features = feature_chain.get().features;
#ifdef __APPLE__
portability_features = feature_chain.get<vk::PhysicalDevicePortabilitySubsetFeaturesKHR>();

9
src/video_core/renderer_vulkan/vk_rasterizer.cpp
View File

@ -532,14 +532,17 @@ void Rasterizer::BindBuffers(const Shader::Info& stage, Shader::Backend::Binding
} else if (desc.buffer_type == Shader::BufferType::ReadConstUbo) {
auto& vk_buffer = buffer_cache.GetStreamBuffer();
const u32 ubo_size = stage.flattened_ud_buf.size() * sizeof(u32);
const u64 offset = vk_buffer.Copy(stage.flattened_ud_buf.data(), ubo_size, instance.UniformMinAlignment());
const u64 offset = vk_buffer.Copy(stage.flattened_ud_buf.data(), ubo_size,
instance.UniformMinAlignment());
buffer_infos.emplace_back(vk_buffer.Handle(), offset, ubo_size);
} else if (desc.buffer_type == Shader::BufferType::SharedMemory) {
// Bind a SSBO to act as shared memory in case of not being able to use a workgroup buffer
// Bind a SSBO to act as shared memory in case of not being able to use a workgroup
// buffer
auto& lds_buffer = buffer_cache.GetStreamBuffer();
const auto& cs_program = liverpool->GetCsRegs();
const auto lds_size = cs_program.SharedMemSize() * cs_program.NumWorkgroups();
const auto [data, offset] = lds_buffer.Map(lds_size, instance.StorageMinAlignment());
const auto [data, offset] =
lds_buffer.Map(lds_size, instance.StorageMinAlignment());
std::memset(data, 0, lds_size);
buffer_infos.emplace_back(lds_buffer.Handle(), offset, lds_size);
} else if (instance.IsNullDescriptorSupported()) {