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

radv: Gather debug info about shader args

Browse source 
Gathers names of shader args using the added macros and stores a list of
arg names with additional information to radv_shader_debug_info.

Example output (pipeline.log, RT prolog):
ARGS:
   0. sgpr const_addr user_data offset=0 size=2 name=ac.ring_offsets
   1. sgpr const_addr user_data offset=2 size=1 name=descriptors[0]
   2. sgpr const_addr user_data offset=3 size=1 name=ac.push_constants
   3. sgpr const_addr user_data offset=4 size=1 name=ac.dynamic_descriptors
   4. sgpr const_addr user_data offset=5 size=1 name=ac.rt.traversal_shader_addr
   5. sgpr const_addr user_data offset=6 size=2 name=ac.rt.sbt_descriptors
   6. sgpr const_addr user_data offset=8 size=2 name=ac.rt.launch_size_addr
   7. sgpr value user_data offset=10 size=1 name=ac.rt.dynamic_callable_stack_base
   8. vgpr value offset=0 size=1 name=ac.local_invocation_ids_packed

Reviewed-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/37829>
This commit is contained in:
Konstantin Seurer 2026-03-31 13:54:17 +02:00 • committed by Marge Bot
parent ff2caf1513
commit 480a94fb16
11 changed files with 166 additions and 58 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
src/amd/vulkan/nir/radv_nir_rt_stage_common.c
View file

@ -257,7 +257,8 @@ radv_nir_return_param_from_type(nir_parameter *param, const glsl_type *type, boo
}
void
radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap)
radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap,
struct radv_shader_debug_info *debug)
{
const struct radv_physical_device *pdev = radv_device_physical(device);
@ -280,7 +281,7 @@ radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage
for (unsigned i = 0; i < 3; i++)
stage->info.cs.uses_block_id[i] = true;
radv_declare_shader_args(device, NULL, &stage->info, MESA_SHADER_COMPUTE, MESA_SHADER_NONE, &stage->args);
radv_declare_shader_args(device, NULL, &stage->info, MESA_SHADER_COMPUTE, MESA_SHADER_NONE, &stage->args, debug);
stage->info.user_sgprs_locs = stage->args.user_sgprs_locs;
b.shader->info.workgroup_size[0] = pdev->rt_wave_size;

3
src/amd/vulkan/nir/radv_nir_rt_stage_common.h
View file

@ -159,6 +159,7 @@ struct radv_nir_rt_traversal_result radv_build_traversal(struct radv_device *dev
struct radv_nir_rt_traversal_params *params,
struct radv_ray_tracing_stage_info *info);
void radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap);
void radv_build_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage, bool uses_descriptor_heap,
struct radv_shader_debug_info *debug);
#endif // MESA_RADV_NIR_RT_STAGE_COMMON_H

4
src/amd/vulkan/radv_debug.c
View file

@ -588,6 +588,10 @@ radv_dump_shader(struct radv_device *device, struct radv_pipeline *pipeline, str
fprintf(f, "NIR:\n%s\n", shader->dbg.nir_string);
}
if (shader->dbg.args_string) {
fprintf(f, "ARGS:\n%s\n", shader->dbg.args_string);
}
fprintf(f, "%s IR:\n%s\n", pdev->use_llvm ? "LLVM" : "ACO", shader->dbg.ir_string);
fprintf(f, "DISASM:\n%s\n", shader->dbg.disasm_string);

1
src/amd/vulkan/radv_pipeline_cache.c
View file

@ -56,6 +56,7 @@ radv_shader_destroy(struct vk_device *_device, struct vk_pipeline_cache_object *
free(shader->dbg.nir_string);
free(shader->dbg.disasm_string);
free(shader->dbg.ir_string);
free(shader->dbg.args_string);
free(shader->dbg.statistics);
free(shader->dbg.debug_info);

2
src/amd/vulkan/radv_pipeline_compute.c
View file

@ -114,7 +114,7 @@ radv_compile_cs(struct radv_device *device, struct radv_shader_stage *cs_stage,
radv_nir_shader_info_pass(device, cs_stage->nir, &cs_stage->layout, &cs_stage->key, NULL, RADV_PIPELINE_COMPUTE,
false, &cs_stage->info);
radv_declare_shader_args(device, NULL, &cs_stage->info, MESA_SHADER_COMPUTE, MESA_SHADER_NONE, &cs_stage->args);
radv_declare_shader_args(device, NULL, &cs_stage->info, MESA_SHADER_COMPUTE, MESA_SHADER_NONE, &cs_stage->args, dbg);
cs_stage->info.user_sgprs_locs = cs_stage->args.user_sgprs_locs;
cs_stage->info.inline_push_constant_mask = cs_stage->args.ac.inline_push_const_mask;

13
src/amd/vulkan/radv_pipeline_graphics.c
View file

@ -2235,14 +2235,15 @@ radv_fill_shader_info(struct radv_device *device, const enum radv_pipeline_type
static void
radv_declare_pipeline_args(struct radv_device *device, struct radv_shader_stage *stages,
const struct radv_graphics_state_key *gfx_state, VkShaderStageFlagBits active_nir_stages)
const struct radv_graphics_state_key *gfx_state, VkShaderStageFlagBits active_nir_stages,
struct radv_shader_debug_info *debug)
{
const struct radv_physical_device *pdev = radv_device_physical(device);
enum amd_gfx_level gfx_level = pdev->info.gfx_level;
if (gfx_level >= GFX9 && stages[MESA_SHADER_TESS_CTRL].nir) {
radv_declare_shader_args(device, gfx_state, &stages[MESA_SHADER_TESS_CTRL].info, MESA_SHADER_TESS_CTRL,
MESA_SHADER_VERTEX, &stages[MESA_SHADER_TESS_CTRL].args);
MESA_SHADER_VERTEX, &stages[MESA_SHADER_TESS_CTRL].args, &debug[MESA_SHADER_TESS_CTRL]);
stages[MESA_SHADER_TESS_CTRL].info.user_sgprs_locs = stages[MESA_SHADER_TESS_CTRL].args.user_sgprs_locs;
stages[MESA_SHADER_TESS_CTRL].info.inline_push_constant_mask =
stages[MESA_SHADER_TESS_CTRL].args.ac.inline_push_const_mask;
@ -2259,7 +2260,7 @@ radv_declare_pipeline_args(struct radv_device *device, struct radv_shader_stage
if (gfx_level >= GFX9 && stages[MESA_SHADER_GEOMETRY].nir) {
mesa_shader_stage pre_stage = stages[MESA_SHADER_TESS_EVAL].nir ? MESA_SHADER_TESS_EVAL : MESA_SHADER_VERTEX;
radv_declare_shader_args(device, gfx_state, &stages[MESA_SHADER_GEOMETRY].info, MESA_SHADER_GEOMETRY, pre_stage,
&stages[MESA_SHADER_GEOMETRY].args);
&stages[MESA_SHADER_GEOMETRY].args, &debug[MESA_SHADER_GEOMETRY]);
stages[MESA_SHADER_GEOMETRY].info.user_sgprs_locs = stages[MESA_SHADER_GEOMETRY].args.user_sgprs_locs;
stages[MESA_SHADER_GEOMETRY].info.inline_push_constant_mask =
stages[MESA_SHADER_GEOMETRY].args.ac.inline_push_const_mask;
@ -2272,7 +2273,7 @@ radv_declare_pipeline_args(struct radv_device *device, struct radv_shader_stage
}
u_foreach_bit (i, active_nir_stages) {
radv_declare_shader_args(device, gfx_state, &stages[i].info, i, MESA_SHADER_NONE, &stages[i].args);
radv_declare_shader_args(device, gfx_state, &stages[i].info, i, MESA_SHADER_NONE, &stages[i].args, &debug[i]);
stages[i].info.user_sgprs_locs = stages[i].args.user_sgprs_locs;
stages[i].info.inline_push_constant_mask = stages[i].args.ac.inline_push_const_mask;
}
@ -2312,7 +2313,7 @@ radv_create_gs_copy_shader(struct radv_device *device, struct vk_pipeline_cache
gs_copy_stage.info.type = RADV_SHADER_TYPE_GS_COPY;
radv_declare_shader_args(device, gfx_state, &gs_copy_stage.info, MESA_SHADER_VERTEX, MESA_SHADER_NONE,
&gs_copy_stage.args);
&gs_copy_stage.args, gs_copy_debug);
gs_copy_stage.info.user_sgprs_locs = gs_copy_stage.args.user_sgprs_locs;
gs_copy_stage.info.inline_push_constant_mask = gs_copy_stage.args.ac.inline_push_const_mask;
@ -2852,7 +2853,7 @@ radv_graphics_shaders_compile(struct radv_device *device, struct vk_pipeline_cac
radv_fill_shader_info(device, RADV_PIPELINE_GRAPHICS, gfx_state, stages, active_nir_stages);
radv_declare_pipeline_args(device, stages, gfx_state, active_nir_stages);
radv_declare_pipeline_args(device, stages, gfx_state, active_nir_stages, debug);
radv_foreach_stage (i, active_nir_stages) {
int64_t stage_start = os_time_get_nano();

7
src/amd/vulkan/radv_pipeline_rt.c
View file

@ -405,7 +405,7 @@ radv_rt_nir_to_asm(struct radv_device *device, struct radv_ray_tracing_pipeline
&stage->info);
/* Declare shader arguments. */
radv_declare_shader_args(device, NULL, &stage->info, stage->stage, MESA_SHADER_NONE, &stage->args);
radv_declare_shader_args(device, NULL, &stage->info, stage->stage, MESA_SHADER_NONE, &stage->args, debug);
stage->info.user_sgprs_locs = stage->args.user_sgprs_locs;
stage->info.inline_push_constant_mask = stage->args.ac.inline_push_const_mask;
@ -1048,7 +1048,8 @@ compile_rt_prolog(struct radv_device *device, struct radv_ray_tracing_pipeline *
uint32_t push_constant_size = 0;
struct radv_shader_stage prolog_stage = {0};
radv_build_rt_prolog(device, &prolog_stage, uses_descriptor_heap);
struct radv_shader_debug_info debug = {0};
radv_build_rt_prolog(device, &prolog_stage, uses_descriptor_heap, &debug);
prolog_stage.nir->options = &pdev->nir_options[MESA_SHADER_COMPUTE];
radv_optimize_nir(prolog_stage.nir, false);
radv_postprocess_nir(device, NULL, &prolog_stage);
@ -1059,7 +1060,7 @@ compile_rt_prolog(struct radv_device *device, struct radv_ray_tracing_pipeline *
NIR_PASS(_, prolog_stage.nir, nir_opt_copy_prop);
NIR_PASS(_, prolog_stage.nir, nir_opt_remove_phis);
pipeline->prolog = radv_compile_rt_prolog(device, &prolog_stage);
pipeline->prolog = radv_compile_rt_prolog(device, &prolog_stage, &debug);
bool has_traversal = !!pipeline->base.base.shaders[MESA_SHADER_INTERSECTION];

12
src/amd/vulkan/radv_shader.c
View file

@ -3454,7 +3454,8 @@ radv_create_trap_handler_shader(struct radv_device *device)
info.type = RADV_SHADER_TYPE_TRAP_HANDLER;
struct radv_shader_args args;
radv_declare_shader_args(device, NULL, &info, stage, MESA_SHADER_NONE, &args);
struct radv_shader_debug_info debug = {};
radv_declare_shader_args(device, NULL, &info, stage, MESA_SHADER_NONE, &args, &debug);
#if AMD_LLVM_AVAILABLE
if (options.dump_shader || options.record_ir)
@ -3474,7 +3475,7 @@ radv_create_trap_handler_shader(struct radv_device *device)
radv_postprocess_binary_config(device, binary, &args);
struct radv_shader *shader;
radv_shader_create_uncached(device, binary, false, NULL, NULL, &shader);
radv_shader_create_uncached(device, binary, false, NULL, &debug, &shader);
radv_parse_binary_debug_info(device, binary, &shader->dbg);
if (options.dump_shader) {
@ -3513,7 +3514,8 @@ radv_aco_build_shader_part(void **bin, uint32_t num_sgprs, uint32_t num_vgprs, c
}
struct radv_shader *
radv_compile_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage)
radv_compile_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage,
struct radv_shader_debug_info *debug)
{
const struct radv_physical_device *pdev = radv_device_physical(device);
struct radv_instance *instance = radv_physical_device_instance(pdev);
@ -3547,7 +3549,7 @@ radv_compile_rt_prolog(struct radv_device *device, struct radv_shader_stage *sta
binary->info = stage->info;
radv_postprocess_binary_config(device, binary, &stage->args);
radv_shader_create_uncached(device, binary, false, NULL, NULL, &prolog);
radv_shader_create_uncached(device, binary, false, NULL, debug, &prolog);
if (!prolog || radv_parse_binary_debug_info(device, binary, &prolog->dbg) != VK_SUCCESS)
goto done;
@ -3588,7 +3590,7 @@ radv_create_vs_prolog(struct radv_device *device, const struct radv_vs_prolog_ke
struct radv_graphics_state_key gfx_state = {0};
radv_declare_shader_args(device, &gfx_state, &info, key->next_stage,
key->next_stage != MESA_SHADER_VERTEX ? MESA_SHADER_VERTEX : MESA_SHADER_NONE, &args);
key->next_stage != MESA_SHADER_VERTEX ? MESA_SHADER_VERTEX : MESA_SHADER_NONE, &args, NULL);
info.user_sgprs_locs = args.user_sgprs_locs;
info.inline_push_constant_mask = args.ac.inline_push_const_mask;

4
src/amd/vulkan/radv_shader.h
View file

@ -428,6 +428,7 @@ struct radv_shader_debug_info {
char *nir_string;
char *disasm_string;
char *ir_string;
char *args_string;
struct amd_stats *statistics;
struct ac_shader_debug_info *debug_info;
uint32_t debug_info_count;
@ -569,7 +570,8 @@ void radv_free_shader_memory(struct radv_device *device, union radv_shader_arena
struct radv_shader *radv_create_trap_handler_shader(struct radv_device *device);
struct radv_shader *radv_compile_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage);
struct radv_shader *radv_compile_rt_prolog(struct radv_device *device, struct radv_shader_stage *stage,
struct radv_shader_debug_info *debug);
struct radv_shader_part *radv_shader_part_create(struct radv_device *device, struct radv_shader_part_binary *binary,
unsigned wave_size);

169
src/amd/vulkan/radv_shader_args.c
View file

@ -14,6 +14,8 @@
#include "radv_physical_device.h"
#include "radv_shader.h"
#include "util/memstream.h"
struct user_sgpr_info {
uint64_t inline_push_constant_mask;
bool inlined_all_push_consts;
@ -50,6 +52,10 @@ allocate_inline_push_consts(const struct radv_shader_info *info, struct user_sgp
struct radv_shader_args_state {
struct radv_shader_args *args;
bool gather_debug_info;
void *ctx;
const char *arg_names[AC_MAX_ARGS];
BITSET_DECLARE(user_data, AC_MAX_ARGS);
};
static void
@ -66,18 +72,44 @@ add_ud_arg(struct radv_shader_args_state *state, unsigned size, enum ac_arg_type
ud_info->num_sgprs += size;
state->args->num_user_sgprs += size;
if (state->gather_debug_info)
BITSET_SET(state->user_data, arg->arg_index);
}
#define RADV_ADD_UD_ARG(state, size, type, arg, ud_index) add_ud_arg(state, size, type, &(state)->args->arg, ud_index)
#define RADV_ADD_UD_ARG(state, size, type, arg, ud_index) \
do { \
add_ud_arg(state, size, type, &(state)->args->arg, ud_index); \
if ((state)->gather_debug_info) { \
(state)->arg_names[(state)->args->arg.arg_index] = #arg; \
} \
} while (false)
#define RADV_ADD_UD_ARRAY_ARG(state, size, type, arg, array_index, ud_index) \
add_ud_arg(state, size, type, &(state)->args->arg[array_index], ud_index)
do { \
add_ud_arg(state, size, type, &(state)->args->arg[array_index], ud_index); \
if ((state)->gather_debug_info) { \
(state)->arg_names[(state)->args->arg[array_index].arg_index] = \
ralloc_asprintf((state)->ctx, "%s[%u]", #arg, array_index); \
} \
} while (false)
#define RADV_ADD_ARG(state, regfile, size, type, arg) \
ac_add_arg(&(state)->args->ac, regfile, size, type, &(state)->args->arg)
do { \
ac_add_arg(&(state)->args->ac, regfile, size, type, &(state)->args->arg); \
if ((state)->gather_debug_info) { \
(state)->arg_names[(state)->args->arg.arg_index] = #arg; \
} \
} while (false)
#define RADV_ADD_ARRAY_ARG(state, regfile, size, type, arg, array_index) \
ac_add_arg(&(state)->args->ac, regfile, size, type, &(state)->args->arg[array_index])
do { \
ac_add_arg(&(state)->args->ac, regfile, size, type, &(state)->args->arg[array_index]); \
if ((state)->gather_debug_info) { \
(state)->arg_names[(state)->args->arg[array_index].arg_index] = \
ralloc_asprintf((state)->ctx, "%s[%u]", #arg, array_index); \
} \
} while (false)
#define RADV_ADD_NULL_ARG(state, regfile, size, type) ac_add_arg(&(state)->args->ac, regfile, size, type, NULL)
@ -881,56 +913,117 @@ declare_shader_args(struct radv_shader_args_state *state, const struct radv_devi
}
}
static void
radv_gather_shader_args_debug_info(struct radv_shader_args_state *state, struct radv_shader_debug_info *debug)
{
char *data = NULL;
size_t size = 0;
struct u_memstream mem;
if (u_memstream_open(&mem, &data, &size)) {
FILE *const memf = u_memstream_get(&mem);
for (uint32_t i = 0; i < state->args->ac.arg_count; i++) {
fprintf(memf, " %u.", i);
switch (state->args->ac.args[i].file) {
case AC_ARG_SGPR:
fprintf(memf, " sgpr");
break;
case AC_ARG_VGPR:
fprintf(memf, " vgpr");
break;
}
switch (state->args->ac.args[i].type) {
case AC_ARG_VALUE:
fprintf(memf, " value");
break;
case AC_ARG_CONST_ADDR:
fprintf(memf, " const_addr");
break;
}
if (state->args->ac.args[i].skip)
fprintf(memf, " skip");
if (state->args->ac.args[i].pending_vmem)
fprintf(memf, " pending_vmem");
if (state->args->ac.args[i].preserved)
fprintf(memf, " preserved");
if (BITSET_TEST(state->user_data, i))
fprintf(memf, " user_data");
fprintf(memf, " offset=%u size=%u name=%s\n", state->args->ac.args[i].offset, state->args->ac.args[i].size,
state->arg_names[i] ? state->arg_names[i] : "(null)");
}
u_memstream_close(&mem);
}
debug->args_string = malloc(size + 1);
if (debug->args_string) {
memcpy(debug->args_string, data, size);
debug->args_string[size] = 0;
}
free(data);
}
void
radv_declare_shader_args(const struct radv_device *device, const struct radv_graphics_state_key *gfx_state,
const struct radv_shader_info *info, mesa_shader_stage stage, mesa_shader_stage previous_stage,
struct radv_shader_args *args)
struct radv_shader_args *args, struct radv_shader_debug_info *debug)
{
struct radv_shader_args_state state = {
.args = args,
};
declare_shader_args(&state, device, gfx_state, info, stage, previous_stage, NULL);
struct user_sgpr_info user_sgpr_info = {};
if (mesa_shader_stage_is_rt(stage))
return;
if (!mesa_shader_stage_is_rt(stage)) {
declare_shader_args(&state, device, gfx_state, info, stage, previous_stage, NULL);
uint32_t num_user_sgprs = args->num_user_sgprs;
if (info->loads_push_constants)
num_user_sgprs++;
if (info->loads_dynamic_offsets) {
num_user_sgprs++;
if (info->loads_dynamic_descriptors_offset_addr)
uint32_t num_user_sgprs = args->num_user_sgprs;
if (info->loads_push_constants)
num_user_sgprs++;
}
const struct radv_physical_device *pdev = radv_device_physical(device);
const enum amd_gfx_level gfx_level = pdev->info.gfx_level;
uint32_t available_sgprs = gfx_level >= GFX9 && stage != MESA_SHADER_COMPUTE && stage != MESA_SHADER_TASK ? 32 : 16;
uint32_t remaining_sgprs = available_sgprs - num_user_sgprs;
struct user_sgpr_info user_sgpr_info = {
.remaining_sgprs = remaining_sgprs,
};
if (info->descriptor_heap) {
assert(user_sgpr_info.remaining_sgprs >= RADV_MAX_HEAPS);
user_sgpr_info.remaining_sgprs -= RADV_MAX_HEAPS;
} else {
const uint32_t num_desc_set = util_bitcount(info->desc_set_used_mask);
if (info->force_indirect_descriptors || remaining_sgprs < num_desc_set) {
user_sgpr_info.indirect_all_descriptor_sets = true;
user_sgpr_info.remaining_sgprs--;
} else {
user_sgpr_info.remaining_sgprs -= num_desc_set;
if (info->loads_dynamic_offsets) {
num_user_sgprs++;
if (info->loads_dynamic_descriptors_offset_addr)
num_user_sgprs++;
}
const struct radv_physical_device *pdev = radv_device_physical(device);
const enum amd_gfx_level gfx_level = pdev->info.gfx_level;
uint32_t available_sgprs =
gfx_level >= GFX9 && stage != MESA_SHADER_COMPUTE && stage != MESA_SHADER_TASK ? 32 : 16;
uint32_t remaining_sgprs = available_sgprs - num_user_sgprs;
user_sgpr_info.remaining_sgprs = remaining_sgprs;
if (info->descriptor_heap) {
assert(user_sgpr_info.remaining_sgprs >= RADV_MAX_HEAPS);
user_sgpr_info.remaining_sgprs -= RADV_MAX_HEAPS;
} else {
const uint32_t num_desc_set = util_bitcount(info->desc_set_used_mask);
if (info->force_indirect_descriptors || remaining_sgprs < num_desc_set) {
user_sgpr_info.indirect_all_descriptor_sets = true;
user_sgpr_info.remaining_sgprs--;
} else {
user_sgpr_info.remaining_sgprs -= num_desc_set;
}
}
if (!info->merged_shader_compiled_separately)
allocate_inline_push_consts(info, &user_sgpr_info);
}
if (!info->merged_shader_compiled_separately)
allocate_inline_push_consts(info, &user_sgpr_info);
state.gather_debug_info = debug && device->keep_shader_info;
if (state.gather_debug_info) {
state.ctx = ralloc_context(NULL);
state.gather_debug_info &= !!state.ctx;
}
declare_shader_args(&state, device, gfx_state, info, stage, previous_stage, &user_sgpr_info);
if (state.gather_debug_info)
radv_gather_shader_args_debug_info(&state, debug);
ralloc_free(state.ctx);
}
void

4
src/amd/vulkan/radv_shader_args.h
View file

@ -136,10 +136,12 @@ struct radv_graphics_state_key;
struct radv_shader_info;
struct radv_ps_epilog_key;
struct radv_device;
struct radv_shader_debug_info;
void radv_declare_shader_args(const struct radv_device *device, const struct radv_graphics_state_key *gfx_state,
const struct radv_shader_info *info, mesa_shader_stage stage,
mesa_shader_stage previous_stage, struct radv_shader_args *args);
mesa_shader_stage previous_stage, struct radv_shader_args *args,
struct radv_shader_debug_info *debug);
void radv_declare_ps_epilog_args(const struct radv_device *device, const struct radv_ps_epilog_key *key,
struct radv_shader_args *args);