Mesa (main): shader_info: Move subgroup_size out of cs and make it an enum

[GitLab Mirror]

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

Author: Jason Ekstrand <jason.ekstrand at collabora.com>
Date:   Thu Jul  7 13:09:22 2022 -0500

shader_info: Move subgroup_size out of cs and make it an enum

Reviewed-by: Caio Oliveira <caio.oliveira at intel.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/17337>

---

 src/compiler/shader_enums.h             | 17 +++++++++++++++++
 src/compiler/shader_info.h              |  7 ++-----
 src/compiler/spirv/spirv_to_nir.c       |  2 +-
 src/intel/compiler/brw_simd_selection.c |  7 ++++---
 4 files changed, 24 insertions(+), 9 deletions(-)

diff --git a/src/compiler/shader_enums.h b/src/compiler/shader_enums.h
index 0d93726d9e6..0ab79323a6b 100644
--- a/src/compiler/shader_enums.h
+++ b/src/compiler/shader_enums.h
@@ -26,6 +26,8 @@
 #ifndef SHADER_ENUMS_H
 #define SHADER_ENUMS_H
 
+#include "util/macros.h"
+
 #include <stdbool.h>
 
 /* Project-wide (GL and Vulkan) maximum. */
@@ -1215,6 +1217,21 @@ enum cl_sampler_filter_mode {
 #define MAT_BIT_FRONT_INDEXES         (1<<MAT_ATTRIB_FRONT_INDEXES)
 #define MAT_BIT_BACK_INDEXES          (1<<MAT_ATTRIB_BACK_INDEXES)
 
+/** An enum representing what kind of input gl_SubgroupSize is. */
+enum PACKED gl_subgroup_size
+{
+   /** Actual subgroup size, whatever that happens to be */
+   SUBGROUP_SIZE_VARYING = 0,
+
+   /* These enums are specifically chosen so that the value of the enum is
+    * also the subgroup size.  If any new values are added, they must respect
+    * this invariant.
+    */
+   SUBGROUP_SIZE_REQUIRE_8   = 8,  /**< VK_EXT_subgroup_size_control */
+   SUBGROUP_SIZE_REQUIRE_16  = 16, /**< VK_EXT_subgroup_size_control */
+   SUBGROUP_SIZE_REQUIRE_32  = 32, /**< VK_EXT_subgroup_size_control */
+};
+
 #ifdef __cplusplus
 } /* extern "C" */
 #endif
diff --git a/src/compiler/shader_info.h b/src/compiler/shader_info.h
index 86d7ca14505..9229c301689 100644
--- a/src/compiler/shader_info.h
+++ b/src/compiler/shader_info.h
@@ -237,6 +237,8 @@ typedef struct shader_info {
     */
    uint16_t workgroup_size[3];
 
+   enum gl_subgroup_size subgroup_size;
+
    /* Transform feedback buffer strides in dwords, max. 1K - 4. */
    uint8_t xfb_stride[MAX_XFB_BUFFERS];
 
@@ -481,11 +483,6 @@ typedef struct shader_info {
           */
          enum gl_derivative_group derivative_group:2;
 
-         /**
-          * Explicit subgroup size if set by the shader, otherwise 0.
-          */
-         unsigned subgroup_size;
-
          /**
           * pointer size is:
           *   AddressingModelLogical:    0    (default)
diff --git a/src/compiler/spirv/spirv_to_nir.c b/src/compiler/spirv/spirv_to_nir.c
index 177acd5dede..bab812fdc12 100644
--- a/src/compiler/spirv/spirv_to_nir.c
+++ b/src/compiler/spirv/spirv_to_nir.c
@@ -5262,7 +5262,7 @@ vtn_handle_execution_mode(struct vtn_builder *b, struct vtn_value *entry_point,
 
    case SpvExecutionModeSubgroupSize:
       vtn_assert(b->shader->info.stage == MESA_SHADER_KERNEL);
-      b->shader->info.cs.subgroup_size = mode->operands[0];
+      b->shader->info.subgroup_size = mode->operands[0];
       break;
 
    case SpvExecutionModeSubgroupUniformControlFlowKHR:
diff --git a/src/intel/compiler/brw_simd_selection.c b/src/intel/compiler/brw_simd_selection.c
index 0c9bf1f4fd2..2ab9bfeed89 100644
--- a/src/intel/compiler/brw_simd_selection.c
+++ b/src/intel/compiler/brw_simd_selection.c
@@ -41,9 +41,10 @@ brw_required_dispatch_width(const struct shader_info *info,
       required = (unsigned)subgroup_size_type;
    }
 
-   if (gl_shader_stage_is_compute(info->stage) && info->cs.subgroup_size > 0) {
-      assert(required == 0 || required == info->cs.subgroup_size);
-      required = info->cs.subgroup_size;
+   if (gl_shader_stage_is_compute(info->stage) &&
+       info->subgroup_size >= SUBGROUP_SIZE_REQUIRE_8) {
+      assert(required == 0 || required == info->subgroup_size);
+      required = info->subgroup_size;
    }
 
    return required;