video_core: Attempt no2 for specialization

Author: raphaelthegreat

CMakeLists.txt

@@ -521,6 +521,8 @@ set(VIDEO_CORE src/video_core/amdgpu/liverpool.cpp
                src/video_core/renderer_vulkan/vk_resource_pool.h
                src/video_core/renderer_vulkan/vk_scheduler.cpp
                src/video_core/renderer_vulkan/vk_scheduler.h
+               src/video_core/renderer_vulkan/vk_shader_cache.cpp
+               src/video_core/renderer_vulkan/vk_shader_cache.h
                src/video_core/renderer_vulkan/vk_shader_util.cpp
                src/video_core/renderer_vulkan/vk_shader_util.h
                src/video_core/renderer_vulkan/vk_swapchain.cpp

src/core/libraries/avplayer/avplayer.cpp

@@ -120,7 +120,6 @@ bool PS4_SYSV_ABI sceAvPlayerGetVideoDataEx(SceAvPlayerHandle handle,
 }
 
 SceAvPlayerHandle PS4_SYSV_ABI sceAvPlayerInit(SceAvPlayerInitData* data) {
-    return nullptr;
     LOG_TRACE(Lib_AvPlayer, "called");
     if (data == nullptr) {
         return nullptr;

src/core/libraries/kernel/thread_management.cpp

@@ -1066,16 +1066,7 @@ ScePthread PThreadPool::Create() {
         }
     }
 
-#ifdef _WIN64
     auto* ret = new PthreadInternal{};
-#else
-    // TODO: Linux specific hack
-    static u8* hint_address = reinterpret_cast<u8*>(0x7FFFFC000ULL);
-    auto* ret = reinterpret_cast<PthreadInternal*>(
-        mmap(hint_address, sizeof(PthreadInternal), PROT_READ | PROT_WRITE,
-             MAP_PRIVATE | MAP_ANONYMOUS | MAP_FIXED, -1, 0));
-    hint_address += Common::AlignUp(sizeof(PthreadInternal), 4_KB);
-#endif
     ret->is_free = false;
     ret->is_detached = false;
     ret->is_almost_done = false;

src/shader_recompiler/backend/spirv/emit_spirv.cpp

@@ -201,6 +201,12 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
     if (info.has_image_query) {
         ctx.AddCapability(spv::Capability::ImageQuery);
     }
+    if (info.uses_lane_id) {
+        ctx.AddCapability(spv::Capability::GroupNonUniform);
+    }
+    if (info.uses_group_quad) {
+        ctx.AddCapability(spv::Capability::GroupNonUniformQuad);
+    }
     switch (program.info.stage) {
     case Stage::Compute: {
         const std::array<u32, 3> workgroup_size{program.info.workgroup_size};
@@ -219,10 +225,6 @@ void DefineEntryPoint(const IR::Program& program, EmitContext& ctx, Id main) {
         } else {
             ctx.AddExecutionMode(main, spv::ExecutionMode::OriginUpperLeft);
         }
-        if (info.uses_group_quad) {
-            ctx.AddCapability(spv::Capability::GroupNonUniform);
-            ctx.AddCapability(spv::Capability::GroupNonUniformQuad);
-        }
         if (info.has_discard) {
             ctx.AddCapability(spv::Capability::DemoteToHelperInvocationEXT);
         }

src/shader_recompiler/backend/spirv/spirv_emit_context.cpp

@@ -132,6 +132,7 @@ const VectorIds& GetAttributeType(EmitContext& ctx, AmdGpu::NumberFormat fmt) {
     case AmdGpu::NumberFormat::SnormNz:
     case AmdGpu::NumberFormat::Sscaled:
     case AmdGpu::NumberFormat::Uscaled:
+    case AmdGpu::NumberFormat::Srgb:
         return ctx.F32;
     case AmdGpu::NumberFormat::Sint:
         return ctx.S32;
@@ -140,7 +141,7 @@ const VectorIds& GetAttributeType(EmitContext& ctx, AmdGpu::NumberFormat fmt) {
     default:
         break;
     }
-    throw InvalidArgument("Invalid attribute type {}", fmt);
+    UNREACHABLE_MSG("Invalid attribute type {}", fmt);
 }
 
 EmitContext::SpirvAttribute EmitContext::GetAttributeInfo(AmdGpu::NumberFormat fmt, Id id) {
@@ -161,7 +162,7 @@ EmitContext::SpirvAttribute EmitContext::GetAttributeInfo(AmdGpu::NumberFormat f
     default:
         break;
     }
-    throw InvalidArgument("Invalid attribute type {}", fmt);
+    UNREACHABLE_MSG("Invalid attribute type {}", fmt);
 }
 
 void EmitContext::DefineBufferOffsets() {
@@ -204,6 +205,11 @@ Id MakeDefaultValue(EmitContext& ctx, u32 default_value) {
 }
 
 void EmitContext::DefineInputs() {
+    if (info.uses_lane_id) {
+        subgroup_local_invocation_id = DefineVariable(
+            U32[1], spv::BuiltIn::SubgroupLocalInvocationId, spv::StorageClass::Input);
+        Decorate(subgroup_local_invocation_id, spv::Decoration::Flat);
+    }
     switch (stage) {
     case Stage::Vertex: {
         vertex_index = DefineVariable(U32[1], spv::BuiltIn::VertexIndex, spv::StorageClass::Input);
@@ -238,9 +244,6 @@ void EmitContext::DefineInputs() {
         break;
     }
     case Stage::Fragment:
-        subgroup_local_invocation_id = DefineVariable(
-            U32[1], spv::BuiltIn::SubgroupLocalInvocationId, spv::StorageClass::Input);
-        Decorate(subgroup_local_invocation_id, spv::Decoration::Flat);
         frag_coord = DefineVariable(F32[4], spv::BuiltIn::FragCoord, spv::StorageClass::Input);
         frag_depth = DefineVariable(F32[1], spv::BuiltIn::FragDepth, spv::StorageClass::Output);
         front_facing = DefineVariable(U1[1], spv::BuiltIn::FrontFacing, spv::StorageClass::Input);
@@ -354,12 +357,12 @@ void EmitContext::DefineBuffers() {
     };
 
     for (const auto& desc : info.buffers) {
-        const auto sharp = desc.GetVsharp(info);
+        const auto sharp = desc.GetSharp(info);
         const bool is_storage = desc.IsStorage(sharp);
         const auto* data_types = True(desc.used_types & IR::Type::F32) ? &F32 : &U32;
         const Id data_type = (*data_types)[1];
         const Id record_array_type{is_storage ? TypeRuntimeArray(data_type)
-                                              : TypeArray(data_type, ConstU32(desc.length))};
+                                              : TypeArray(data_type, ConstU32(sharp.NumDwords()))};
         const Id struct_type{define_struct(record_array_type, desc.is_instance_data)};
 
         const auto storage_class =
@@ -369,6 +372,9 @@ void EmitContext::DefineBuffers() {
         const Id id{AddGlobalVariable(struct_pointer_type, storage_class)};
         Decorate(id, spv::Decoration::Binding, binding);
         Decorate(id, spv::Decoration::DescriptorSet, 0U);
+        if (is_storage && !desc.is_written) {
+            Decorate(id, spv::Decoration::NonWritable);
+        }
         Name(id, fmt::format("{}_{}", is_storage ? "ssbo" : "cbuf", desc.sgpr_base));
 
         buffers.push_back({
@@ -503,17 +509,8 @@ Id ImageType(EmitContext& ctx, const ImageResource& desc, Id sampled_type) {
 
 void EmitContext::DefineImagesAndSamplers() {
     for (const auto& image_desc : info.images) {
-        const VectorIds* data_types = [&] {
-            switch (image_desc.nfmt) {
-            case AmdGpu::NumberFormat::Uint:
-                return &U32;
-            case AmdGpu::NumberFormat::Sint:
-                return &S32;
-            default:
-                return &F32;
-            }
-        }();
-        const Id sampled_type = data_types->Get(1);
+        const VectorIds& data_types = GetAttributeType(*this, image_desc.nfmt);
+        const Id sampled_type = data_types[1];
         const Id image_type{ImageType(*this, image_desc, sampled_type)};
         const Id pointer_type{TypePointer(spv::StorageClass::UniformConstant, image_type)};
         const Id id{AddGlobalVariable(pointer_type, spv::StorageClass::UniformConstant)};
@@ -522,7 +519,7 @@ void EmitContext::DefineImagesAndSamplers() {
         Name(id, fmt::format("{}_{}{}_{:02x}", stage, "img", image_desc.sgpr_base,
                              image_desc.dword_offset));
         images.push_back({
-            .data_types = data_types,
+            .data_types = &data_types,
             .id = id,
             .sampled_type = image_desc.is_storage ? sampled_type : TypeSampledImage(image_type),
             .pointer_type = pointer_type,
@@ -531,13 +528,12 @@ void EmitContext::DefineImagesAndSamplers() {
         interfaces.push_back(id);
         ++binding;
     }
-
-    image_u32 = TypePointer(spv::StorageClass::Image, U32[1]);
-
+    if (std::ranges::any_of(info.images, &ImageResource::is_atomic)) {
+        image_u32 = TypePointer(spv::StorageClass::Image, U32[1]);
+    }
     if (info.samplers.empty()) {
         return;
     }
-
     sampler_type = TypeSampler();
     sampler_pointer_type = TypePointer(spv::StorageClass::UniformConstant, sampler_type);
     for (const auto& samp_desc : info.samplers) {
@@ -553,7 +549,7 @@ void EmitContext::DefineImagesAndSamplers() {
 }
 
 void EmitContext::DefineSharedMemory() {
-    static constexpr size_t DefaultSharedMemSize = 16_KB;
+    static constexpr size_t DefaultSharedMemSize = 2_KB;
     if (!info.uses_shared) {
         return;
     }

src/shader_recompiler/frontend/translate/export.cpp

@@ -1,14 +1,12 @@
 // SPDX-FileCopyrightText: Copyright 2024 shadPS4 Emulator Project
 // SPDX-License-Identifier: GPL-2.0-or-later
 
-#include "common/logging/log.h"
 #include "shader_recompiler/frontend/translate/translate.h"
 
 namespace Shader::Gcn {
 
 void Translator::EmitExport(const GcnInst& inst) {
     if (ir.block->has_multiple_predecessors && info.stage == Stage::Fragment) {
-        LOG_WARNING(Render_Recompiler, "An ambiguous export appeared in translation");
         ir.Discard(ir.LogicalNot(ir.GetExec()));
     }
 

src/shader_recompiler/frontend/translate/translate.cpp

@@ -399,7 +399,6 @@ void Translator::EmitFetch(const GcnInst& inst) {
             info.buffers.push_back({
                 .sgpr_base = attrib.sgpr_base,
                 .dword_offset = attrib.dword_offset,
-                .length = buffer.num_records,
                 .used_types = IR::Type::F32,
                 .is_instance_data = true,
             });

src/shader_recompiler/frontend/translate/vector_alu.cpp

@@ -415,14 +415,20 @@ void Translator::V_ADDC_U32(const GcnInst& inst) {
     const auto src0 = GetSrc<IR::U32>(inst.src[0]);
     const auto src1 = GetSrc<IR::U32>(inst.src[1]);
 
-    IR::U32 scarry;
+    IR::U1 carry;
     if (inst.src_count == 3) { // VOP3
-        IR::U1 thread_bit{ir.GetThreadBitScalarReg(IR::ScalarReg(inst.src[2].code))};
-        scarry = IR::U32{ir.Select(thread_bit, ir.Imm32(1), ir.Imm32(0))};
+        if (inst.src[2].field == OperandField::VccLo) {
+            carry = ir.GetVcc();
+        } else if (inst.src[2].field == OperandField::ScalarGPR) {
+            carry = ir.GetThreadBitScalarReg(IR::ScalarReg(inst.src[2].code));
+        } else {
+            UNREACHABLE();
+        }
     } else { // VOP2
-        scarry = ir.GetVccLo();
+        carry = ir.GetVcc();
     }
 
+    const IR::U32 scarry = IR::U32{ir.Select(carry, ir.Imm32(1), ir.Imm32(0))};
     const IR::U32 result = ir.IAdd(ir.IAdd(src0, src1), scarry);
 
     const IR::VectorReg dst_reg{inst.dst[0].code};

src/shader_recompiler/ir/passes/resource_tracking_pass.cpp

@@ -3,6 +3,7 @@
 
 #include <algorithm>
 #include <boost/container/small_vector.hpp>
+#include "common/alignment.h"
 #include "shader_recompiler/ir/basic_block.h"
 #include "shader_recompiler/ir/breadth_first_search.h"
 #include "shader_recompiler/ir/ir_emitter.h"
@@ -195,7 +196,6 @@ class Descriptors {
                    desc.inline_cbuf == existing.inline_cbuf;
         })};
         auto& buffer = buffer_resources[index];
-        ASSERT(buffer.length == desc.length);
         buffer.used_types |= desc.used_types;
         buffer.is_written |= desc.is_written;
         return index;
@@ -227,7 +227,7 @@ class Descriptors {
                 return true;
             }
             // Samplers with different bindings might still be the same.
-            return existing.GetSsharp(info) == desc.GetSsharp(info);
+            return existing.GetSharp(info) == desc.GetSharp(info);
         })};
         return index;
     }
@@ -342,19 +342,6 @@ SharpLocation TrackSharp(const IR::Inst* inst) {
     };
 }
 
-static u32 BufferLength(const AmdGpu::Buffer& buffer) {
-    const auto stride = buffer.GetStride();
-    if (stride < sizeof(f32)) {
-        ASSERT(sizeof(f32) % stride == 0);
-        return (((buffer.num_records - 1) / sizeof(f32)) + 1) * stride;
-    } else if (stride == sizeof(f32)) {
-        return buffer.num_records;
-    } else {
-        ASSERT(stride % sizeof(f32) == 0);
-        return buffer.num_records * (stride / sizeof(f32));
-    }
-}
-
 s32 TryHandleInlineCbuf(IR::Inst& inst, Info& info, Descriptors& descriptors,
                         AmdGpu::Buffer& cbuf) {
 
@@ -381,7 +368,6 @@ s32 TryHandleInlineCbuf(IR::Inst& inst, Info& info, Descriptors& descriptors,
     return descriptors.Add(BufferResource{
         .sgpr_base = std::numeric_limits<u32>::max(),
         .dword_offset = 0,
-        .length = BufferLength(cbuf),
         .used_types = BufferDataType(inst, cbuf.GetNumberFmt()),
         .inline_cbuf = cbuf,
     });
@@ -399,7 +385,6 @@ void PatchBufferInstruction(IR::Block& block, IR::Inst& inst, Info& info,
         binding = descriptors.Add(BufferResource{
             .sgpr_base = sharp.sgpr_base,
             .dword_offset = sharp.dword_offset,
-            .length = BufferLength(buffer),
             .used_types = BufferDataType(inst, buffer.GetNumberFmt()),
             .is_written = IsBufferStore(inst),
         });

src/shader_recompiler/ir/passes/shader_info_collection_pass.cpp

@@ -50,6 +50,9 @@ void Visit(Info& info, IR::Inst& inst) {
     case IR::Opcode::ImageQueryLod:
         info.has_image_query = true;
         break;
+    case IR::Opcode::LaneId:
+        info.uses_lane_id = true;
+        break;
     default:
         break;
     }