fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-05-06 02:58:05 +02:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

radeonsi: implement get_compute_state_info

Browse source 
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>
This commit is contained in:
Karol Herbst 2022-11-18 23:26:01 +01:00 • committed by Marge Bot
parent 0c730f98c5
commit e5ef95e31f
1 changed files with 19 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

19
src/gallium/drivers/radeonsi/si_compute.c
View file

@ -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;