Mesa (master): spirv: Handle most execution modes earlier
GitLab Mirror
gitlab-mirror at kemper.freedesktop.org
Mon Jul 27 17:30:48 UTC 2020
Module: Mesa
Branch: master
Commit: 12dd5455f43c8fe24f50c78a880270cf8cc023c5
URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=12dd5455f43c8fe24f50c78a880270cf8cc023c5
Author: Caio Marcelo de Oliveira Filho <caio.oliveira at intel.com>
Date: Mon Jul 6 22:58:25 2020 -0700
spirv: Handle most execution modes earlier
For convenience in e68871f6a44 ("spirv: Handle constants and types
before execution modes") we moved all execution mode parsing after the
constants and types, so that those using OpExecutionModeId could be
handled together.
Later in 84781e1f1d8 ("spirv/nir: keep track of SPV_KHR_float_controls
execution modes") we had to parse certain non-ID execution modes
before handling constants.
Instead of handling just the float controls related execution modes
early, handle all modes that don't need an ID. This is a more
"natural" split and will allow other type handling to rely on
execution mode in the future.
Reviewed-by: Jason Ekstrand <jason at jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6062>
---
src/compiler/spirv/spirv_to_nir.c | 134 +++++++++++++++++++++-----------------
1 file changed, 73 insertions(+), 61 deletions(-)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 799926eb037..96c0c0767db 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -4441,14 +4441,7 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
b->shader->info.cs.local_size[2] = mode->operands[2];
break;
- case SpvExecutionModeLocalSizeId:
- b->shader->info.cs.local_size[0] = vtn_constant_uint(b, mode->operands[0]);
- b->shader->info.cs.local_size[1] = vtn_constant_uint(b, mode->operands[1]);
- b->shader->info.cs.local_size[2] = vtn_constant_uint(b, mode->operands[2]);
- break;
-
case SpvExecutionModeLocalSizeHint:
- case SpvExecutionModeLocalSizeHintId:
break; /* Nothing to do with this */
case SpvExecutionModeOutputVertices:
@@ -4578,8 +4571,60 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
case SpvExecutionModeDenormFlushToZero:
case SpvExecutionModeSignedZeroInfNanPreserve:
case SpvExecutionModeRoundingModeRTE:
- case SpvExecutionModeRoundingModeRTZ:
- /* Already handled in vtn_handle_rounding_mode_in_execution_mode() */
+ case SpvExecutionModeRoundingModeRTZ: {
+ unsigned execution_mode = 0;
+ switch (mode->exec_mode) {
+ case SpvExecutionModeDenormPreserve:
+ switch (mode->operands[0]) {
+ case 16: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP16; break;
+ case 32: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP32; break;
+ case 64: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP64; break;
+ default: vtn_fail("Floating point type not supported");
+ }
+ break;
+ case SpvExecutionModeDenormFlushToZero:
+ switch (mode->operands[0]) {
+ case 16: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP16; break;
+ case 32: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP32; break;
+ case 64: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP64; break;
+ default: vtn_fail("Floating point type not supported");
+ }
+ break;
+ case SpvExecutionModeSignedZeroInfNanPreserve:
+ switch (mode->operands[0]) {
+ case 16: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP16; break;
+ case 32: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP32; break;
+ case 64: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP64; break;
+ default: vtn_fail("Floating point type not supported");
+ }
+ break;
+ case SpvExecutionModeRoundingModeRTE:
+ switch (mode->operands[0]) {
+ case 16: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP16; break;
+ case 32: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP32; break;
+ case 64: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP64; break;
+ default: vtn_fail("Floating point type not supported");
+ }
+ break;
+ case SpvExecutionModeRoundingModeRTZ:
+ switch (mode->operands[0]) {
+ case 16: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP16; break;
+ case 32: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP32; break;
+ case 64: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP64; break;
+ default: vtn_fail("Floating point type not supported");
+ }
+ break;
+ default:
+ break;
+ }
+
+ b->shader->info.float_controls_execution_mode |= execution_mode;
+ break;
+ }
+
+ case SpvExecutionModeLocalSizeId:
+ case SpvExecutionModeLocalSizeHintId:
+ /* Handled later by vtn_handle_execution_mode_id(). */
break;
default:
@@ -4590,60 +4635,28 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
}
static void
-vtn_handle_rounding_mode_in_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
- const struct vtn_decoration *mode, void *data)
+vtn_handle_execution_mode_id(struct vtn_builder *b, struct vtn_value *entry_point,
+ const struct vtn_decoration *mode, UNUSED void *data)
{
- vtn_assert(b->entry_point == entry_point);
- unsigned execution_mode = 0;
+ vtn_assert(b->entry_point == entry_point);
- switch(mode->exec_mode) {
- case SpvExecutionModeDenormPreserve:
- switch (mode->operands[0]) {
- case 16: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP16; break;
- case 32: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP32; break;
- case 64: execution_mode = FLOAT_CONTROLS_DENORM_PRESERVE_FP64; break;
- default: vtn_fail("Floating point type not supported");
- }
- break;
- case SpvExecutionModeDenormFlushToZero:
- switch (mode->operands[0]) {
- case 16: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP16; break;
- case 32: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP32; break;
- case 64: execution_mode = FLOAT_CONTROLS_DENORM_FLUSH_TO_ZERO_FP64; break;
- default: vtn_fail("Floating point type not supported");
- }
- break;
- case SpvExecutionModeSignedZeroInfNanPreserve:
- switch (mode->operands[0]) {
- case 16: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP16; break;
- case 32: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP32; break;
- case 64: execution_mode = FLOAT_CONTROLS_SIGNED_ZERO_INF_NAN_PRESERVE_FP64; break;
- default: vtn_fail("Floating point type not supported");
- }
- break;
- case SpvExecutionModeRoundingModeRTE:
- switch (mode->operands[0]) {
- case 16: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP16; break;
- case 32: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP32; break;
- case 64: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTE_FP64; break;
- default: vtn_fail("Floating point type not supported");
- }
+ switch (mode->exec_mode) {
+ case SpvExecutionModeLocalSizeId:
+ b->shader->info.cs.local_size[0] = vtn_constant_uint(b, mode->operands[0]);
+ b->shader->info.cs.local_size[1] = vtn_constant_uint(b, mode->operands[1]);
+ b->shader->info.cs.local_size[2] = vtn_constant_uint(b, mode->operands[2]);
break;
- case SpvExecutionModeRoundingModeRTZ:
- switch (mode->operands[0]) {
- case 16: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP16; break;
- case 32: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP32; break;
- case 64: execution_mode = FLOAT_CONTROLS_ROUNDING_MODE_RTZ_FP64; break;
- default: vtn_fail("Floating point type not supported");
- }
+
+ case SpvExecutionModeLocalSizeHintId:
+ /* Nothing to do with this hint. */
break;
default:
+ /* Nothing to do. Literal execution modes already handled by
+ * vtn_handle_execution_mode(). */
break;
}
-
- b->shader->info.float_controls_execution_mode |= execution_mode;
}
static bool
@@ -5438,12 +5451,9 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
if (stage == MESA_SHADER_GEOMETRY)
b->shader->info.gs.invocations = 1;
- /* Parse rounding mode execution modes. This has to happen earlier than
- * other changes in the execution modes since they can affect, for example,
- * the result of the floating point constants.
- */
+ /* Parse execution modes. */
vtn_foreach_execution_mode(b, b->entry_point,
- vtn_handle_rounding_mode_in_execution_mode, NULL);
+ vtn_handle_execution_mode, NULL);
b->specializations = spec;
b->num_specializations = num_spec;
@@ -5452,9 +5462,11 @@ spirv_to_nir(const uint32_t *words, size_t word_count,
words = vtn_foreach_instruction(b, words, word_end,
vtn_handle_variable_or_type_instruction);
- /* Parse execution modes */
+ /* Parse execution modes that depend on IDs. Must happen after we have
+ * constants parsed.
+ */
vtn_foreach_execution_mode(b, b->entry_point,
- vtn_handle_execution_mode, NULL);
+ vtn_handle_execution_mode_id, NULL);
if (b->workgroup_size_builtin) {
vtn_assert(b->workgroup_size_builtin->type->type ==
More information about the mesa-commit
mailing list