shader_recompiler: don't compile cs with higher shared memory than supported

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);

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;
 

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;

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: