From 6ef8413799bcbcd39760ebd343fda2e196552c68 Mon Sep 17 00:00:00 2001
From: polyproxy <47796739+polybiusproxy@users.noreply.github.com>
Date: Fri, 17 Jan 2025 21:13:23 +0100
Subject: [PATCH] don't compile cs with higher shared memory than supported

---
 src/shader_recompiler/backend/spirv/spirv_emit_context.cpp | 4 ++++
 src/shader_recompiler/runtime_info.h                       | 1 +
 src/video_core/renderer_vulkan/vk_instance.h               | 5 +++++
 src/video_core/renderer_vulkan/vk_pipeline_cache.cpp       | 1 +
 4 files changed, 11 insertions(+)

diff --git a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
index 7d492a38444..295bef75d8b 100644
--- a/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
+++ b/src/shader_recompiler/backend/spirv/spirv_emit_context.cpp
@@ -847,6 +847,10 @@ void EmitContext::DefineSharedMemory() {
     if (shared_memory_size == 0) {
         shared_memory_size = DefaultSharedMemSize;
     }
+
+    const u32 max_shared_memory_size = runtime_info.cs_info.max_shared_memory_size;
+    ASSERT(shared_memory_size <= max_shared_memory_size);
+
     const u32 num_elements{Common::DivCeil(shared_memory_size, 4U)};
     const Id type{TypeArray(U32[1], ConstU32(num_elements))};
     shared_memory_u32_type = TypePointer(spv::StorageClass::Workgroup, type);
diff --git a/src/shader_recompiler/runtime_info.h b/src/shader_recompiler/runtime_info.h
index cf49b08793c..8a5e0a1a6b9 100644
--- a/src/shader_recompiler/runtime_info.h
+++ b/src/shader_recompiler/runtime_info.h
@@ -198,6 +198,7 @@ struct FragmentRuntimeInfo {
 
 struct ComputeRuntimeInfo {
     u32 shared_memory_size;
+    u32 max_shared_memory_size;
     std::array<u32, 3> workgroup_size;
     std::array<bool, 3> tgid_enable;
 
diff --git a/src/video_core/renderer_vulkan/vk_instance.h b/src/video_core/renderer_vulkan/vk_instance.h
index e0d4d0b4d81..1b0c276dcaf 100644
--- a/src/video_core/renderer_vulkan/vk_instance.h
+++ b/src/video_core/renderer_vulkan/vk_instance.h
@@ -239,6 +239,11 @@ class Instance {
         return subgroup_size;
     }
 
+    /// Returns the maximum size of compute shared memory.
+    u32 MaxComputeSharedMemorySize() const {
+        return properties.limits.maxComputeSharedMemorySize;
+    }
+
     /// Returns the maximum supported elements in a texel buffer
     u32 MaxTexelBufferElements() const {
         return properties.limits.maxTexelBufferElements;
diff --git a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
index f934943899f..86f6dcc7a8a 100644
--- a/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
+++ b/src/video_core/renderer_vulkan/vk_pipeline_cache.cpp
@@ -183,6 +183,7 @@ const Shader::RuntimeInfo& PipelineCache::BuildRuntimeInfo(Stage stage, LogicalS
         info.cs_info.tgid_enable = {cs_pgm.IsTgidEnabled(0), cs_pgm.IsTgidEnabled(1),
                                     cs_pgm.IsTgidEnabled(2)};
         info.cs_info.shared_memory_size = cs_pgm.SharedMemSize();
+        info.cs_info.max_shared_memory_size = instance.MaxComputeSharedMemorySize();
         break;
     }
     default: