Mesa (main): nv50/ir/nir: fix smem size for GL

[GitLab Mirror]

Module: Mesa
Branch: main
Commit: ff55412f40a7d588e47e2c35d175aea0ac3cfe95
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=ff55412f40a7d588e47e2c35d175aea0ac3cfe95

Author: Karol Herbst <kherbst at redhat.com>
Date:   Thu May 27 18:59:37 2021 +0200

nv50/ir/nir: fix smem size for GL

Originally I fixed the case where the nir itself has a shared mem size of
0, but the frontend (e.g. clover) set it to some other value.

But st/mesa sets the shared mem size on the state object as well and we
end up actually doubling the value in the driver as we set smemSize to the
value from the state object before calling into the compiler.

So just max the value instead.

Fixes the compute_shader.shared-max CTS test.

Fixes: dc667b1f192 ("nv50/ir/nir: fix smem size")
Signed-off-by: Karol Herbst <kherbst at redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11047>

---

 src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp | 2 +-
 1 file changed, 1 insertion(+), 1 deletion(-)

diff --git a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
index d6dcc6b11af..75e3ef14786 100644
--- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
+++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp
@@ -1301,7 +1301,7 @@ Converter::parseNIR()
       info->prop.cp.numThreads[0] = nir->info.workgroup_size[0];
       info->prop.cp.numThreads[1] = nir->info.workgroup_size[1];
       info->prop.cp.numThreads[2] = nir->info.workgroup_size[2];
-      info_out->bin.smemSize += nir->info.shared_size;
+      info_out->bin.smemSize = std::max(info_out->bin.smemSize, nir->info.shared_size);
       break;
    case Program::TYPE_FRAGMENT:
       info_out->prop.fp.earlyFragTests = nir->info.fs.early_fragment_tests;