Mesa (master): spirv: Ignore WorkgroupSize in non-compute stages
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Thu Mar 11 20:51:16 UTC 2021
Module: Mesa
Branch: master
Commit: 4a408ff7ea96ca792c5b0e589fd85bf4490f3973
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=4a408ff7ea96ca792c5b0e589fd85bf4490f3973
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>
---
src/compiler/spirv/spirv_to_nir.c | 6 +++++-
1 file changed, 5 insertions(+), 1 deletion(-)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index a765fd5cbf1..0316cbd7099 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -2175,7 +2175,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
@@ -5922,6 +5925,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));
More information about the mesa-commit
mailing list