From 8e00c1956d7df4754115badf344f42b2150e7214 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Thu, 27 May 2021 18:59:37 +0200 Subject: [PATCH] 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 Part-of: (cherry picked from commit ff55412f40a7d588e47e2c35d175aea0ac3cfe95) --- .pick_status.json | 2 +- src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/.pick_status.json b/.pick_status.json index a6da7e38ed5..31eaa463839 100644 --- a/.pick_status.json +++ b/.pick_status.json @@ -697,7 +697,7 @@ "description": "nv50/ir/nir: fix smem size for GL", "nominated": true, "nomination_type": 1, - "resolution": 0, + "resolution": 1, "main_sha": null, "because_sha": "dc667b1f192d33d073832a50b0e920734f9fb8ef" }, 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 59c0ed89eb6..7aaa8cbb40b 100644 --- a/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp +++ b/src/gallium/drivers/nouveau/codegen/nv50_ir_from_nir.cpp @@ -1292,7 +1292,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;