From 67998c4a027c2446a93ca5c7a85fe9e20825b460 Mon Sep 17 00:00:00 2001 From: IndecisiveTurtle <47210458+raphaelthegreat@users.noreply.github.com> Date: Fri, 14 Feb 2025 12:16:34 +0200 Subject: [PATCH] clang format --- .../spirv/emit_spirv_context_get_set.cpp | 12 +++---- .../backend/spirv/spirv_emit_context.cpp | 35 ++++++++++++------- .../backend/spirv/spirv_emit_context.h | 4 +-- .../frontend/translate/data_share.cpp | 6 ++-- .../ir/passes/resource_tracking_pass.cpp | 3 +- .../ir/passes/shared_memory_barrier_pass.cpp | 15 +++++--- .../passes/shared_memory_to_storage_pass.cpp | 24 ++++++++----- .../renderer_vulkan/vk_instance.cpp | 14 ++++---- .../renderer_vulkan/vk_rasterizer.cpp | 9 +++-- 9 files changed, 75 insertions(+), 47 deletions(-) diff --git a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp index 04ac28d24..cc7b7e097 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv_context_get_set.cpp @@ -165,9 +165,11 @@ using BufferAlias = EmitContext::BufferAlias; Id EmitReadConst(EmitContext& ctx, IR::Inst* inst) { const u32 flatbuf_off_dw = inst->Flags(); 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 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)); } } diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp index 50d846e45..8946f7d82 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -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]); diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.h b/src/shader_recompiler/backend/spirv/spirv_emit_context.h index 29d883bcf..75605fe91 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.h +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.h @@ -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); diff --git a/src/shader_recompiler/frontend/translate/data_share.cpp b/src/shader_recompiler/frontend/translate/data_share.cpp index 36ed9eaad..460f8913c 100644 --- a/src/shader_recompiler/frontend/translate/data_share.cpp +++ b/src/shader_recompiler/frontend/translate/data_share.cpp @@ -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; } diff --git a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp index bdcb65fc3..c5bfe5796 100644 --- a/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp +++ b/src/shader_recompiler/ir/passes/resource_tracking_pass.cpp @@ -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]; diff --git a/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp b/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp index 3a83c601e..0ee52cf19 100644 --- a/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp +++ b/src/shader_recompiler/ir/passes/shared_memory_barrier_pass.cpp @@ -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; diff --git a/src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp b/src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp index 517852fae..25aaf257c 100644 --- a/src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp +++ b/src/shader_recompiler/ir/passes/shared_memory_to_storage_pass.cpp @@ -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: diff --git a/src/video_core/renderer_vulkan/vk_instance.cpp b/src/video_core/renderer_vulkan/vk_instance.cpp index 761ef6fff..a17f8c9c2 100644 --- a/src/video_core/renderer_vulkan/vk_instance.cpp +++ b/src/video_core/renderer_vulkan/vk_instance.cpp @@ -203,14 +203,12 @@ std::string Instance::GetDriverVersionName() { } bool Instance::CreateDevice() { - const vk::StructureChain feature_chain = - physical_device - .getFeatures2(); + 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(); diff --git a/src/video_core/renderer_vulkan/vk_rasterizer.cpp b/src/video_core/renderer_vulkan/vk_rasterizer.cpp index a9ce4b4e2..19e5863a7 100644 --- a/src/video_core/renderer_vulkan/vk_rasterizer.cpp +++ b/src/video_core/renderer_vulkan/vk_rasterizer.cpp @@ -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()) {