radeonsi: implement get_compute_state_info

Signed-off-by: Karol Herbst <kherbst@redhat.com>
Reviewed-by: Pierre-Eric Pelloux-Prayer <pierre-eric.pelloux-prayer@amd.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/19232>

diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c
index cb1f117..9ed4f90 100644 (file)
--- a/src/gallium/drivers/radeonsi/si_compute.c
+++ b/src/gallium/drivers/radeonsi/si_compute.c
@@ -29,6 +29,7 @@
 #include "amd_kernel_code_t.h"
 #include "nir/tgsi_to_nir.h"
 #include "si_build_pm4.h"
+#include "si_shader_internal.h"
 #include "util/u_async_debug.h"
 #include "util/u_memory.h"
 #include "util/u_upload_mgr.h"
@@ -296,6 +297,23 @@ static void *si_create_compute_state(struct pipe_context *ctx, const struct pipe
    return program;
 }
 
+static void si_get_compute_state_info(struct pipe_context *ctx, void *state,
+                                      struct pipe_compute_state_object_info *info)
+{
+   struct si_compute *program = (struct si_compute *)state;
+   struct si_shader_selector *sel = &program->sel;
+
+   assert(program->ir_type != PIPE_SHADER_IR_NATIVE);
+
+   /* Wait because we need the compilation to finish first */
+   util_queue_fence_wait(&sel->ready);
+
+   uint8_t wave_size = program->shader.wave_size;
+   info->private_memory = DIV_ROUND_UP(program->shader.config.scratch_bytes_per_wave, wave_size);
+   info->preferred_simd_size = wave_size;
+   info->max_threads = si_get_max_workgroup_size(&program->shader);
+}
+
 static void si_bind_compute_state(struct pipe_context *ctx, void *state)
 {
    struct si_context *sctx = (struct si_context *)ctx;
@@ -1042,6 +1060,7 @@ void si_init_compute_functions(struct si_context *sctx)
    sctx->b.create_compute_state = si_create_compute_state;
    sctx->b.delete_compute_state = si_delete_compute_state;
    sctx->b.bind_compute_state = si_bind_compute_state;
+   sctx->b.get_compute_state_info = si_get_compute_state_info;
    sctx->b.set_compute_resources = si_set_compute_resources;
    sctx->b.set_global_binding = si_set_global_binding;
    sctx->b.launch_grid = si_launch_grid;