From b72d9509392174caf88f6e90b3a2725e047706f4 Mon Sep 17 00:00:00 2001 From: Karol Herbst Date: Sun, 4 Jun 2023 16:35:44 +0200 Subject: [PATCH] gallium: add simd_sizes to pipe_compute_state_object_info Signed-off-by: Karol Herbst Reviewed-by: Alyssa Rosenzweig Part-of: --- src/gallium/drivers/iris/iris_program.c | 1 + src/gallium/drivers/llvmpipe/lp_state_cs.c | 3 ++- src/gallium/drivers/nouveau/nv50/nv50_state.c | 1 + src/gallium/drivers/nouveau/nvc0/nvc0_state.c | 1 + src/gallium/drivers/panfrost/pan_shader.c | 3 ++- src/gallium/drivers/radeonsi/si_compute.c | 1 + src/gallium/include/pipe/p_state.h | 5 +++++ 7 files changed, 13 insertions(+), 2 deletions(-) diff --git a/src/gallium/drivers/iris/iris_program.c b/src/gallium/drivers/iris/iris_program.c index b86bc27..f4b8cfa 100644 --- a/src/gallium/drivers/iris/iris_program.c +++ b/src/gallium/drivers/iris/iris_program.c @@ -2505,6 +2505,7 @@ iris_get_compute_state_info(struct pipe_context *ctx, void *state, info->max_threads = MIN2(1024, 32 * screen->devinfo->max_cs_workgroup_threads); info->private_memory = 0; info->preferred_simd_size = 32; + info->simd_sizes = 8 | 16 | 32; list_for_each_entry_safe(struct iris_compiled_shader, shader, &ish->variants, link) { diff --git a/src/gallium/drivers/llvmpipe/lp_state_cs.c b/src/gallium/drivers/llvmpipe/lp_state_cs.c index cdb29f0..97c094c 100644 --- a/src/gallium/drivers/llvmpipe/lp_state_cs.c +++ b/src/gallium/drivers/llvmpipe/lp_state_cs.c @@ -922,7 +922,8 @@ llvmpipe_get_compute_state_info(struct pipe_context *pipe, void *cs, struct nir_shader* nir = shader->base.ir.nir; info->max_threads = 1024; - info->preferred_simd_size = lp_native_vector_width / 32; + info->simd_sizes = lp_native_vector_width / 32; + info->preferred_simd_size = info->simd_sizes; // TODO: this is a bad estimate, but not much we can do without actually compiling the shaders info->private_memory = nir->scratch_size; } diff --git a/src/gallium/drivers/nouveau/nv50/nv50_state.c b/src/gallium/drivers/nouveau/nv50/nv50_state.c index 57170a5..3791240 100644 --- a/src/gallium/drivers/nouveau/nv50/nv50_state.c +++ b/src/gallium/drivers/nouveau/nv50/nv50_state.c @@ -894,6 +894,7 @@ nv50_get_compute_state_info(struct pipe_context *pipe, void *hwcso, info->max_threads = MIN2(ROUND_DOWN_TO(threads, 32), 512); info->private_memory = prog->tls_space; info->preferred_simd_size = 32; + info->simd_sizes = 32; } static void diff --git a/src/gallium/drivers/nouveau/nvc0/nvc0_state.c b/src/gallium/drivers/nouveau/nvc0/nvc0_state.c index 08a2db0..4f79546 100644 --- a/src/gallium/drivers/nouveau/nvc0/nvc0_state.c +++ b/src/gallium/drivers/nouveau/nvc0/nvc0_state.c @@ -811,6 +811,7 @@ nvc0_get_compute_state_info(struct pipe_context *pipe, void *hwcso, info->max_threads = MIN2(ROUND_DOWN_TO(threads, 32), 1024); info->private_memory = prog->hdr[1] & 0xfffff0; info->preferred_simd_size = 32; + info->simd_sizes = 32; } static void diff --git a/src/gallium/drivers/panfrost/pan_shader.c b/src/gallium/drivers/panfrost/pan_shader.c index 33c864e..42bb2f7 100644 --- a/src/gallium/drivers/panfrost/pan_shader.c +++ b/src/gallium/drivers/panfrost/pan_shader.c @@ -499,7 +499,8 @@ panfrost_get_compute_state_info(struct pipe_context *pipe, void *cso, info->max_threads = panfrost_max_thread_count(dev->arch, cs->info.work_reg_count); info->private_memory = cs->info.tls_size; - info->preferred_simd_size = pan_subgroup_size(dev->arch); + info->simd_sizes = pan_subgroup_size(dev->arch); + info->preferred_simd_size = info->simd_sizes; } void diff --git a/src/gallium/drivers/radeonsi/si_compute.c b/src/gallium/drivers/radeonsi/si_compute.c index fcd4242..28aaa9d 100644 --- a/src/gallium/drivers/radeonsi/si_compute.c +++ b/src/gallium/drivers/radeonsi/si_compute.c @@ -294,6 +294,7 @@ static void si_get_compute_state_info(struct pipe_context *ctx, void *state, 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->simd_sizes = wave_size; info->max_threads = si_get_max_workgroup_size(&program->shader); } diff --git a/src/gallium/include/pipe/p_state.h b/src/gallium/include/pipe/p_state.h index 1c3f77f..549e4d2 100644 --- a/src/gallium/include/pipe/p_state.h +++ b/src/gallium/include/pipe/p_state.h @@ -1062,6 +1062,11 @@ struct pipe_compute_state_object_info unsigned preferred_simd_size; /** + * Bitmask of supported SIMD sizes. + */ + unsigned simd_sizes; + + /** * How much private memory does this CSO require per thread (a.k.a. NIR scratch memory). */ unsigned private_memory; -- 2.7.4