Mesa (staging/21.0): spirv: Ignore WorkgroupSize in non-compute stages

[GitLab Mirror]

Module: Mesa
Branch: staging/21.0
Commit: 0f84de9cbb12dda2663f21447a8700c6264b493d
URL:    http://cgit.freedesktop.org/mesa/mesa/commit/?id=0f84de9cbb12dda2663f21447a8700c6264b493d

Author: Pierre Moreau <dev at pmoreau.org>
Date:   Thu Mar  4 22:52:09 2021 +0100

spirv: Ignore WorkgroupSize in non-compute stages

If a SPIR-V module contains for example both a geometry and a compute
shader, when processing the geometry shader its vertices out, input
primitive and output primitive attributes would get overwritten by the
value of the WorkgroupSize.

```
; SPIR-V
; Version: 1.5
; Generator: Khronos; 17
; Bound: 12
; Schema: 0
               OpCapability Geometry
               OpCapability Shader
          %1 = OpExtInstImport "GLSL.std.450"
               OpMemoryModel Logical GLSL450
               OpEntryPoint Geometry %main "main"
               OpEntryPoint GLCompute %main_0 "main"
               OpExecutionMode %main InputPoints
               OpExecutionMode %main Invocations 1
               OpExecutionMode %main OutputTriangleStrip
               OpExecutionMode %main OutputVertices 4
               OpExecutionMode %main_0 LocalSize 1 1 1
               OpSource GLSL 460
               OpSource GLSL 460
               OpName %main "main"
               OpName %main_0 "main"
               OpModuleProcessed "Linked by SPIR-V Tools Linker"
               OpDecorate %gl_WorkGroupSize BuiltIn WorkgroupSize
       %void = OpTypeVoid
          %6 = OpTypeFunction %void
       %uint = OpTypeInt 32 0
     %v3uint = OpTypeVector %uint 3
     %uint_1 = OpConstant %uint 1
%gl_WorkGroupSize = OpConstantComposite %v3uint %uint_1 %uint_1 %uint_1
       %main = OpFunction %void None %6
         %10 = OpLabel
               OpReturn
               OpFunctionEnd
     %main_0 = OpFunction %void None %6
         %11 = OpLabel
               OpReturn
               OpFunctionEnd
```

Running spirv_to_nir on the SPIR-V sample above and for the geometry
entry point would say that (among others):

* vertices out: 1
* input primitive: LINES
* output primitive: LINES

By removing any reference to `%gl_WorkGroupSize`, the output would
change to (among others):

* vertices out: 4
* input primitive: POINTS
* output primitive: TRIANGLE_STRIP

Fixes: 7d862ef5302 ("spirv: Rework handling of spec constant workgroup size built-ins")

v2:
* Move the check from inside `handle_workgroup_size_decoration_cb()` to
  its caller (Caio Marcelo de Oliveira Filho )
* Add an assert on the shader stage before using
  `workgroup_size_builtin` (Caio Marcelo de Oliveira Filho )

Reviewed-by: Caio Marcelo de Oliveira Filho <caio.oliveira at intel.com>
Signed-off-by: Pierre Moreau <dev at pmoreau.org>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/9418>
(cherry picked from commit 4a408ff7ea96ca792c5b0e589fd85bf4490f3973)

---

 .pick_status.json                 | 2 +-
 src/compiler/spirv/spirv_to_nir.c | 6 +++++-
 2 files changed, 6 insertions(+), 2 deletions(-)

diff --git a/.pick_status.json b/.pick_status.json
index fee329f5955..4baa2af27eb 100644
--- a/.pick_status.json
+++ b/.pick_status.json
@@ -1849,7 +1849,7 @@
         "description": "spirv: Ignore WorkgroupSize in non-compute stages",
         "nominated": true,
         "nomination_type": 1,
-        "resolution": 0,
+        "resolution": 1,
         "master_sha": null,
         "because_sha": "7d862ef53024b70458929c1665af43573b4d07ff"
     },
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 6864bb61cb2..6a784812ffd 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -2166,7 +2166,10 @@ vtn_handle_constant(struct vtn_builder *b, SpvOp opcode,
    }
 
    /* Now that we have the value, update the workgroup size if needed */
-   vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb, NULL);
+   if (b->entry_point_stage == MESA_SHADER_COMPUTE ||
+       b->entry_point_stage == MESA_SHADER_KERNEL)
+      vtn_foreach_decoration(b, val, handle_workgroup_size_decoration_cb,
+                             NULL);
 }
 
 static void
@@ -5897,6 +5900,7 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
                                  vtn_handle_execution_mode_id, NULL);
 
    if (b->workgroup_size_builtin) {
+      vtn_assert(stage == MESA_SHADER_COMPUTE || stage == MESA_SHADER_KERNEL);
       vtn_assert(b->workgroup_size_builtin->type->type ==
                  glsl_vector_type(GLSL_TYPE_UINT, 3));