From 64337f004d9249c4408fec75bd1bbcc0f2a1408d Mon Sep 17 00:00:00 2001
From: ameerj <52414509+ameerj@users.noreply.github.com>
Date: Thu, 20 May 2021 23:38:38 -0400
Subject: glsl: Fix "reg" allocing

based on glasm with some tweaks
---
 .../backend/glsl/emit_context.cpp                  | 28 ++++++++++++++++++----
 1 file changed, 24 insertions(+), 4 deletions(-)

(limited to 'src/shader_recompiler/backend/glsl/emit_context.cpp')

diff --git a/src/shader_recompiler/backend/glsl/emit_context.cpp b/src/shader_recompiler/backend/glsl/emit_context.cpp
index e2a9885f0..b0e46cc07 100644
--- a/src/shader_recompiler/backend/glsl/emit_context.cpp
+++ b/src/shader_recompiler/backend/glsl/emit_context.cpp
@@ -11,19 +11,39 @@ namespace Shader::Backend::GLSL {
 EmitContext::EmitContext(IR::Program& program, [[maybe_unused]] Bindings& bindings,
                          const Profile& profile_)
     : info{program.info}, profile{profile_} {
-    std::string header = "#version 450 core\n";
-    header += "layout(local_size_x=1, local_size_y=1, local_size_z=1) in;";
+    std::string header = "#version 450\n";
+    if (program.stage == Stage::Compute) {
+        header += fmt::format("layout(local_size_x={},local_size_y={},local_size_z={}) in;\n",
+                              program.workgroup_size[0], program.workgroup_size[1],
+                              program.workgroup_size[2]);
+    }
     code += header;
     DefineConstantBuffers();
-    code += "void main(){";
+    DefineStorageBuffers();
+    code += "void main(){\n";
 }
 
 void EmitContext::DefineConstantBuffers() {
     if (info.constant_buffer_descriptors.empty()) {
         return;
     }
+    u32 binding{};
     for (const auto& desc : info.constant_buffer_descriptors) {
-        Add("uniform uint c{}[{}];", desc.index, desc.count);
+        Add("layout(std140,binding={}) uniform cbuf_{}{{uint cbuf{}[];}};", binding, binding,
+            desc.index, desc.count);
+        ++binding;
+    }
+}
+
+void EmitContext::DefineStorageBuffers() {
+    if (info.storage_buffers_descriptors.empty()) {
+        return;
+    }
+    u32 binding{};
+    for (const auto& desc : info.storage_buffers_descriptors) {
+        Add("layout(std430,binding={}) buffer buff_{}{{uint buff{}[];}};", binding, binding,
+            desc.cbuf_index, desc.count);
+        ++binding;
     }
 }
 
-- 
cgit v1.2.3