From 9e9481699d18f7a56c92d492b5f231af1cb2ab54 Mon Sep 17 00:00:00 2001
From: Jordan Justen <jordan.l.justen@intel.com>
Date: Fri, 2 Nov 2018 01:56:52 -0700
Subject: [PATCH] intel/blorp: Add blorp_get_cs_local_y, blorp_set_cs_dims

Based the blorp_params, blorp_get_cs_local_y returns a recommended
local_y size for a compute shader.

blorp_set_cs_dims sets the compute program dims based on a given
local_y size.

Reworks:
 * Add blorp_set_cs_dims (s-b Jason)

Signed-off-by: Jordan Justen <jordan.l.justen@intel.com>
Reviewed-by: Kenneth Graunke <kenneth@whitecape.org>
Reviewed-by: Jason Ekstrand <jason@jlekstrand.net>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/11564>
---
 src/intel/blorp/blorp_priv.h | 23 +++++++++++++++++++++++
 1 file changed, 23 insertions(+)

diff --git a/src/intel/blorp/blorp_priv.h b/src/intel/blorp/blorp_priv.h
index 95247b75d08..ad1ed80bcc8 100644
--- a/src/intel/blorp/blorp_priv.h
+++ b/src/intel/blorp/blorp_priv.h
@@ -420,6 +420,29 @@ bool
 blorp_ensure_sf_program(struct blorp_batch *batch,
                         struct blorp_params *params);
 
+static inline uint8_t
+blorp_get_cs_local_y(struct blorp_params *params)
+{
+   uint32_t height = params->y1 - params->y0;
+   uint32_t or_ys = params->y0 | params->y1;
+   if (height > 32 || (or_ys & 3) == 0) {
+      return 4;
+   } else if ((or_ys & 1) == 0) {
+      return 2;
+   } else {
+      return 1;
+   }
+}
+
+static inline void
+blorp_set_cs_dims(struct nir_shader *nir, uint8_t local_y)
+{
+   assert(local_y != 0 && (16 % local_y == 0));
+   nir->info.workgroup_size[0] = 16 / local_y;
+   nir->info.workgroup_size[1] = local_y;
+   nir->info.workgroup_size[2] = 1;
+}
+
 const unsigned *
 blorp_compile_cs(struct blorp_context *blorp, void *mem_ctx,
                  struct nir_shader *nir,