Module: Mesa Branch: main Commit: dc3e69af1af861871b2944a389a3027be15a89a7 URL: http://cgit.freedesktop.org/mesa/mesa/commit/?id=dc3e69af1af861871b2944a389a3027be15a89a7
Author: Lionel Landwerlin <lionel.g.landwer...@intel.com> Date: Thu Nov 30 19:13:45 2023 +0200 nir/serialize: untangle printf serialization from a particular stage This allows any stage to carry printf instructions. Signed-off-by: Lionel Landwerlin <lionel.g.landwer...@intel.com> Reviewed-by: Karol Herbst <kher...@redhat.com> Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/26505> --- src/compiler/nir/nir_serialize.c | 4 ++-- src/compiler/shader_info.h | 5 +++++ src/compiler/spirv/vtn_opencl.c | 2 ++ 3 files changed, 9 insertions(+), 2 deletions(-) diff --git a/src/compiler/nir/nir_serialize.c b/src/compiler/nir/nir_serialize.c index cf53cd894e2..64965b847ec 100644 --- a/src/compiler/nir/nir_serialize.c +++ b/src/compiler/nir/nir_serialize.c @@ -2024,7 +2024,7 @@ nir_serialize(struct blob *blob, const nir_shader *nir, bool strip) write_xfb_info(&ctx, nir->xfb_info); - if (nir->info.stage == MESA_SHADER_KERNEL) + if (nir->info.uses_printf) nir_serialize_printf_info(blob, nir->printf_info, nir->printf_info_count); blob_overwrite_uint32(blob, idx_size_offset, ctx.next_idx); @@ -2084,7 +2084,7 @@ nir_deserialize(void *mem_ctx, ctx.nir->xfb_info = read_xfb_info(&ctx); - if (ctx.nir->info.stage == MESA_SHADER_KERNEL) { + if (ctx.nir->info.uses_printf) { ctx.nir->printf_info = nir_deserialize_printf_info(ctx.nir, blob, &ctx.nir->printf_info_count); diff --git a/src/compiler/shader_info.h b/src/compiler/shader_info.h index aee93be814f..a2bb006ea8f 100644 --- a/src/compiler/shader_info.h +++ b/src/compiler/shader_info.h @@ -341,6 +341,11 @@ typedef struct shader_info { */ bool workgroup_size_variable:1; + /** + * Whether the shader uses printf instructions. + */ + bool uses_printf:1; + /** * Set if this shader uses legacy (DX9 or ARB assembly) math rules. * diff --git a/src/compiler/spirv/vtn_opencl.c b/src/compiler/spirv/vtn_opencl.c index d36b851d0eb..f49130dda6c 100644 --- a/src/compiler/spirv/vtn_opencl.c +++ b/src/compiler/spirv/vtn_opencl.c @@ -830,6 +830,8 @@ handle_printf(struct vtn_builder *b, uint32_t opcode, nir_def *fmt_idx = nir_imm_int(&b->nb, info_idx); nir_def *ret = nir_printf(&b->nb, fmt_idx, &deref_var->def); vtn_push_nir_ssa(b, w_dest[1], ret); + + b->nb.shader->info.uses_printf = true; } static nir_def *
