1 files changed, 15 insertions, 18 deletions

diff --git a/src/video_core/renderer_opengl/gl_arb_decompiler.cpp b/src/video_core/renderer_opengl/gl_arb_decompiler.cpp
index 4489abf61..b7e9ed2e9 100644
--- a/src/video_core/renderer_opengl/gl_arb_decompiler.cpp
+++ b/src/video_core/renderer_opengl/gl_arb_decompiler.cpp
@@ -913,11 +913,19 @@ void ARBDecompiler::DeclareCompute() {
     const ComputeInfo& info = registry.GetComputeInfo();
     AddLine("GROUP_SIZE {} {} {};", info.workgroup_size[0], info.workgroup_size[1],
             info.workgroup_size[2]);
-    if (info.shared_memory_size_in_words > 0) {
-        const u32 size_in_bytes = info.shared_memory_size_in_words * 4;
-        AddLine("SHARED_MEMORY {};", size_in_bytes);
-        AddLine("SHARED shared_mem[] = {{program.sharedmem}};");
+    if (info.shared_memory_size_in_words == 0) {
+        return;
+    }
+    const u32 limit = device.GetMaxComputeSharedMemorySize();
+    u32 size_in_bytes = info.shared_memory_size_in_words * 4;
+    if (size_in_bytes > limit) {
+        LOG_ERROR(Render_OpenGL, "Shared memory size {} is clamped to host's limit {}",
+                  size_in_bytes, limit);
+        size_in_bytes = limit;
     }
+
+    AddLine("SHARED_MEMORY {};", size_in_bytes);
+    AddLine("SHARED shared_mem[] = {{program.sharedmem}};");
 }
 
 void ARBDecompiler::DeclareInputAttributes() {
@@ -1283,13 +1291,6 @@ std::string ARBDecompiler::Visit(const Node& node) {
             return "{0, 0, 0, 0}.x";
         }
 
-        const auto buffer_index = [this, &abuf]() -> std::string {
-            if (stage != ShaderType::Geometry) {
-                return "";
-            }
-            return fmt::format("[{}]", Visit(abuf->GetBuffer()));
-        };
-
         const Attribute::Index index = abuf->GetIndex();
         const u32 element = abuf->GetElement();
         const char swizzle = Swizzle(element);
@@ -1395,7 +1396,7 @@ std::string ARBDecompiler::Visit(const Node& node) {
         return {};
     }
 
-    if (const auto cmt = std::get_if<CommentNode>(&*node)) {
+    if ([[maybe_unused]] const auto cmt = std::get_if<CommentNode>(&*node)) {
         // Uncommenting this will generate invalid code. GLASM lacks comments.
         // AddLine("// {}", cmt->GetText());
         return {};
@@ -1703,7 +1704,7 @@ std::string ARBDecompiler::HCastFloat(Operation operation) {
 }
 
 std::string ARBDecompiler::HUnpack(Operation operation) {
-    const std::string operand = Visit(operation[0]);
+    std::string operand = Visit(operation[0]);
     switch (std::get<Tegra::Shader::HalfType>(operation.GetMeta())) {
     case Tegra::Shader::HalfType::H0_H1:
         return operand;
@@ -2053,7 +2054,7 @@ std::string ARBDecompiler::InvocationId(Operation) {
 
 std::string ARBDecompiler::YNegate(Operation) {
     LOG_WARNING(Render_OpenGL, "(STUBBED)");
-    const std::string temporary = AllocTemporary();
+    std::string temporary = AllocTemporary();
     AddLine("MOV.F {}, 1;", temporary);
     return temporary;
 }
@@ -2076,10 +2077,6 @@ std::string ARBDecompiler::ShuffleIndexed(Operation operation) {
 }
 
 std::string ARBDecompiler::Barrier(Operation) {
-    if (!ir.IsDecompiled()) {
-        LOG_ERROR(Render_OpenGL, "BAR used but shader is not decompiled");
-        return {};
-    }
     AddLine("BAR;");
     return {};
 }