nir/range_analysis: Set higher default maximum for max_workgroup_count

Fixes: c2a81ebe19f ("nir: Add default unsigned upper bound configuration.")
Closes: #7676
Reviewed-by: Karol Herbst <kherbst@redhat.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19835>
(cherry picked from commit 2ba55ec504f8391775622c3753ddb03bdcd85aff)

diff --git a/.pick_status.json b/.pick_status.json
index 58d24d5..b4e29e8 100644 (file)
--- a/.pick_status.json
+++ b/.pick_status.json
@@ -1669,7 +1669,7 @@
         "description": "nir/range_analysis: Set higher default maximum for max_workgroup_count",
         "nominated": true,
         "nomination_type": 1,
-        "resolution": 0,
+        "resolution": 1,
         "main_sha": null,
         "because_sha": "c2a81ebe19f98b025b296fcadc279b4358d37345"
     },

diff --git a/src/compiler/nir/nir.h b/src/compiler/nir/nir.h
index 349c8a3..0cf9dfb 100644 (file)
--- a/src/compiler/nir/nir.h
+++ b/src/compiler/nir/nir.h
@@ -5683,6 +5683,7 @@ nir_variable_is_in_block(const nir_variable *var)
    return nir_variable_is_in_ubo(var) || nir_variable_is_in_ssbo(var);
 }
 
+/* See default_ub_config in nir_range_analysis.c for documentation. */
 typedef struct nir_unsigned_upper_bound_config {
    unsigned min_subgroup_size;
    unsigned max_subgroup_size;

diff --git a/src/compiler/nir/nir_range_analysis.c b/src/compiler/nir/nir_range_analysis.c
index 700159b..56fd3f0 100644 (file)
--- a/src/compiler/nir/nir_range_analysis.c
+++ b/src/compiler/nir/nir_range_analysis.c
@@ -1294,8 +1294,19 @@ static const nir_unsigned_upper_bound_config default_ub_config = {
    .min_subgroup_size = 1u,
    .max_subgroup_size = UINT16_MAX,
    .max_workgroup_invocations = UINT16_MAX,
-   .max_workgroup_count = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
+
+   /* max_workgroup_count represents the maximum compute shader / kernel
+    * dispatchable work size. On most hardware, this is essentially
+    * unbounded. On some hardware max_workgroup_count[1] and
+    * max_workgroup_count[2] may be smaller.
+    */
+   .max_workgroup_count = {UINT32_MAX, UINT32_MAX, UINT32_MAX},
+
+   /* max_workgroup_size is the local invocation maximum. This is generally
+    * small the OpenGL 4.2 minimum maximum is 1024.
+    */
    .max_workgroup_size = {UINT16_MAX, UINT16_MAX, UINT16_MAX},
+
    .vertex_attrib_max = {
       UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX,
       UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX, UINT32_MAX,