From a36407bc5e2000956a7b20c727fa18bd38117fe1 Mon Sep 17 00:00:00 2001 From: IndecisiveTurtle <47210458+raphaelthegreat@users.noreply.github.com> Date: Fri, 5 Jul 2024 02:12:26 +0300 Subject: [PATCH] shader_recompiler: Make most features optional --- .../backend/spirv/emit_spirv.cpp | 26 +++++++++++++------ .../backend/spirv/spirv_emit_context.cpp | 16 +++++------- .../ir/passes/shader_info_collection_pass.cpp | 18 +++++++++++++ src/shader_recompiler/runtime_info.h | 5 ++++ src/video_core/amdgpu/liverpool.h | 11 ++++---- .../renderer_vulkan/vk_instance.cpp | 8 +++--- 6 files changed, 58 insertions(+), 26 deletions(-) diff --git a/src/shader_recompiler/backend/spirv/emit_spirv.cpp b/src/shader_recompiler/backend/spirv/emit_spirv.cpp index b1a0845fb..561014a33 100644 --- a/src/shader_recompiler/backend/spirv/emit_spirv.cpp +++ b/src/shader_recompiler/backend/spirv/emit_spirv.cpp @@ -174,14 +174,18 @@ Id DefineMain(EmitContext& ctx, IR::Program& program) { } void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { + const auto& info = program.info; const std::span interfaces(ctx.interfaces.data(), ctx.interfaces.size()); spv::ExecutionModel execution_model{}; ctx.AddCapability(spv::Capability::Image1D); ctx.AddCapability(spv::Capability::Sampled1D); - ctx.AddCapability(spv::Capability::Float16); - ctx.AddCapability(spv::Capability::Int16); - ctx.AddCapability(spv::Capability::StorageImageWriteWithoutFormat); - ctx.AddCapability(spv::Capability::StorageImageExtendedFormats); + if (info.uses_fp16) { + ctx.AddCapability(spv::Capability::Float16); + ctx.AddCapability(spv::Capability::Int16); + } + if (info.has_storage_images) { + ctx.AddCapability(spv::Capability::StorageImageExtendedFormats); + } switch (program.info.stage) { case Stage::Compute: { const std::array workgroup_size{program.info.workgroup_size}; @@ -200,13 +204,19 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) { } else { ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft); } - if (program.info.uses_group_quad) { + if (info.uses_group_quad) { ctx.AddCapability(spv::Capability::GroupNonUniform); ctx.AddCapability(spv::Capability::GroupNonUniformQuad); } - ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); - ctx.AddCapability(spv::Capability::ImageGatherExtended); - ctx.AddCapability(spv::Capability::ImageQuery); + if (info.has_discard) { + ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT); + } + if (info.has_image_gather) { + ctx.AddCapability(spv::Capability::ImageGatherExtended); + } + if (info.has_image_query) { + ctx.AddCapability(spv::Capability::ImageQuery); + } // if (program.info.stores_frag_depth) { // ctx.AddExecutionMode(main, spv::ExecutionMode::DepthReplacing); // } diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp index b4a67ebfb..6d9b25470 100644 --- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp +++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp @@ -74,21 +74,19 @@ Id EmitContext::Def(const IR::Value& value) { void EmitContext::DefineArithmeticTypes() { void_id = Name(TypeVoid(), "void_id"); U1[1] = Name(TypeBool(), "bool_id"); - F16[1] = Name(TypeFloat(16), "f16_id"); + if (info.uses_fp16) { + F16[1] = Name(TypeFloat(16), "f16_id"); + U16 = Name(TypeUInt(16), "u16_id"); + } F32[1] = Name(TypeFloat(32), "f32_id"); - // F64[1] = Name(TypeFloat(64), "f64_id"); S32[1] = Name(TypeSInt(32), "i32_id"); U32[1] = Name(TypeUInt(32), "u32_id"); - // U8 = Name(TypeSInt(8), "u8"); - // S8 = Name(TypeUInt(8), "s8"); - U16 = Name(TypeUInt(16), "u16_id"); - // S16 = Name(TypeSInt(16), "s16_id"); - // U64 = Name(TypeUInt(64), "u64_id"); for (u32 i = 2; i <= 4; i++) { - F16[i] = Name(TypeVector(F16[1], i), fmt::format("f16vec{}_id", i)); + if (info.uses_fp16) { + F16[i] = Name(TypeVector(F16[1], i), fmt::format("f16vec{}_id", i)); + } F32[i] = Name(TypeVector(F32[1], i), fmt::format("f32vec{}_id", i)); - // F64[i] = Name(TypeVector(F64[1], i), fmt::format("f64vec{}_id", i)); S32[i] = Name(TypeVector(S32[1], i), fmt::format("i32vec{}_id", i)); U32[i] = Name(TypeVector(U32[1], i), fmt::format("u32vec{}_id", i)); U1[i] = Name(TypeVector(U1[1], i), fmt::format("bvec{}_id", i)); diff --git a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp index 25d8b9377..1cec237fc 100644 --- a/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp +++ b/src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp @@ -26,9 +26,27 @@ void Visit(Info& info, IR::Inst& inst) { case IR::Opcode::WriteSharedU16: info.uses_shared_u16 = true; break; + case IR::Opcode::ConvertF32F16: + case IR::Opcode::BitCastF16U16: + info.uses_fp16 = true; + break; + case IR::Opcode::ImageWrite: + info.has_storage_images = true; + break; case IR::Opcode::QuadShuffle: info.uses_group_quad = true; break; + case IR::Opcode::Discard: + info.has_discard = true; + break; + case IR::Opcode::ImageGather: + case IR::Opcode::ImageGatherDref: + info.has_image_gather = true; + break; + case IR::Opcode::ImageQueryDimensions: + case IR::Opcode::ImageQueryLod: + info.has_image_query = true; + break; default: break; } diff --git a/src/shader_recompiler/runtime_info.h b/src/shader_recompiler/runtime_info.h index 993207eb4..66d32d4d7 100644 --- a/src/shader_recompiler/runtime_info.h +++ b/src/shader_recompiler/runtime_info.h @@ -169,9 +169,14 @@ struct Info { uintptr_t pgm_base{}; u64 pgm_hash{}; u32 shared_memory_size{}; + bool has_storage_images{}; + bool has_discard{}; + bool has_image_gather{}; + bool has_image_query{}; bool uses_group_quad{}; bool uses_shared_u8{}; bool uses_shared_u16{}; + bool uses_fp16{}; bool translation_failed{}; // indicates that shader has unsupported instructions template diff --git a/src/video_core/amdgpu/liverpool.h b/src/video_core/amdgpu/liverpool.h index 6e19f55da..d18482f6e 100644 --- a/src/video_core/amdgpu/liverpool.h +++ b/src/video_core/amdgpu/liverpool.h @@ -3,12 +3,6 @@ #pragma once -#include "common/assert.h" -#include "common/bit_field.h" -#include "common/types.h" -#include "resource.h" -#include "video_core/amdgpu/pixel_format.h" - #include #include #include @@ -16,6 +10,11 @@ #include #include #include +#include "common/assert.h" +#include "common/bit_field.h" +#include "common/types.h" +#include "video_core/amdgpu/pixel_format.h" +#include "video_core/amdgpu/resource.h" namespace Vulkan { class Rasterizer; diff --git a/src/video_core/renderer_vulkan/vk_instance.cpp b/src/video_core/renderer_vulkan/vk_instance.cpp index 56cafb9d0..3e2135ce7 100644 --- a/src/video_core/renderer_vulkan/vk_instance.cpp +++ b/src/video_core/renderer_vulkan/vk_instance.cpp @@ -114,6 +114,7 @@ bool Instance::CreateDevice() { vk::PhysicalDeviceExtendedDynamicState3FeaturesEXT, vk::PhysicalDeviceCustomBorderColorFeaturesEXT, vk::PhysicalDeviceColorWriteEnableFeaturesEXT, vk::PhysicalDeviceVulkan12Features, + vk::PhysicalDeviceVulkan13Features, vk::PhysicalDeviceWorkgroupMemoryExplicitLayoutFeaturesKHR, vk::PhysicalDeviceDepthClipControlFeaturesEXT>(); const vk::StructureChain properties_chain = @@ -189,6 +190,7 @@ bool Instance::CreateDevice() { }; const auto vk12_features = feature_chain.get(); + const auto vk13_features = feature_chain.get(); vk::StructureChain device_chain = { vk::DeviceCreateInfo{ .queueCreateInfoCount = 1u, @@ -223,9 +225,9 @@ bool Instance::CreateDevice() { .timelineSemaphore = vk12_features.timelineSemaphore, }, vk::PhysicalDeviceVulkan13Features{ - .shaderDemoteToHelperInvocation = true, - .dynamicRendering = true, - .maintenance4 = true, + .shaderDemoteToHelperInvocation = vk13_features.shaderDemoteToHelperInvocation, + .dynamicRendering = vk13_features.dynamicRendering, + .maintenance4 = vk13_features.maintenance4, }, vk::PhysicalDeviceCustomBorderColorFeaturesEXT{ .customBorderColors = true,