fdo-mirrors/mesa
1
0
Fork 0
mirror of https://gitlab.freedesktop.org/mesa/mesa.git synced 2026-01-04 20:00:11 +01:00
Code Issues Projects Releases Packages Wiki Activity Actions 1

aco: add a helper for building a trap handler shader

Browse source 
It's way easier to write a trap handler shader using ACO IR
instead of writing disassembly by hand + clrxasm + copy&paste.

This trap handler is quite simple for now, it just loads a
buffer descriptor from the TMA BO, it saves ttmp0-1 which
contain various info about the faulty instruction, and it
stores some hw registers about the wave/trap status.

Signed-off-by: Samuel Pitoiset <samuel.pitoiset@gmail.com>
Reviewed-by: Daniel Schürmann <daniel@schuermann.dev>
Part-of: <https://gitlab.freedesktop.org/mesa/mesa/-/merge_requests/6384>
This commit is contained in:
Samuel Pitoiset 2020-08-18 18:39:20 +02:00 • committed by Marge Bot
parent a6146aa598
commit 9c46e6fca3
2 changed files with 63 additions and 0 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

60
src/amd/compiler/aco_instruction_selection.cpp
View file

@ -11104,4 +11104,64 @@ void select_gs_copy_shader(Program *program, struct nir_shader *gs_shader,
cleanup_cfg(program);
}
void select_trap_handler_shader(Program *program, struct nir_shader *shader,
ac_shader_config* config,
struct radv_shader_args *args)
{
assert(args->options->chip_class == GFX8);
init_program(program, compute_cs, args->shader_info,
args->options->chip_class, args->options->family, config);
isel_context ctx = {};
ctx.program = program;
ctx.args = args;
ctx.options = args->options;
ctx.stage = program->stage;
ctx.block = ctx.program->create_and_insert_block();
ctx.block->loop_nest_depth = 0;
ctx.block->kind = block_kind_top_level;
program->workgroup_size = 1; /* XXX */
add_startpgm(&ctx);
append_logical_start(ctx.block);
Builder bld(ctx.program, ctx.block);
/* Load the buffer descriptor from TMA. */
bld.smem(aco_opcode::s_load_dwordx4, Definition(PhysReg{ttmp4}, s4),
Operand(PhysReg{tma}, s2), Operand(0u));
/* Store TTMP0-TTMP1. */
bld.smem(aco_opcode::s_buffer_store_dwordx2, Operand(PhysReg{ttmp4}, s4),
Operand(0u), Operand(PhysReg{ttmp0}, s2), memory_sync_info(), true);
uint32_t hw_regs_idx[] = {
2, /* HW_REG_STATUS */
3, /* HW_REG_TRAP_STS */
4, /* HW_REG_HW_ID */
7, /* HW_REG_IB_STS */
};
/* Store some hardware registers. */
for (unsigned i = 0; i < ARRAY_SIZE(hw_regs_idx); i++) {
/* "((size - 1) << 11) | register" */
bld.sopk(aco_opcode::s_getreg_b32, Definition(PhysReg{ttmp8}, s1),
((20 - 1) << 11) | hw_regs_idx[i]);
bld.smem(aco_opcode::s_buffer_store_dword, Operand(PhysReg{ttmp4}, s4),
Operand(8u + i * 4), Operand(PhysReg{ttmp8}, s1), memory_sync_info(), true);
}
program->config->float_mode = program->blocks[0].fp_mode.val;
append_logical_end(ctx.block);
ctx.block->kind |= block_kind_uniform;
bld.sopp(aco_opcode::s_endpgm);
cleanup_cfg(program);
}
}

3
src/amd/compiler/aco_ir.h
View file

@ -1665,6 +1665,9 @@ void select_program(Program *program,
void select_gs_copy_shader(Program *program, struct nir_shader *gs_shader,
ac_shader_config* config,
struct radv_shader_args *args);
void select_trap_handler_shader(Program *program, struct nir_shader *shader,
ac_shader_config* config,
struct radv_shader_args *args);
void lower_wqm(Program* program, live& live_vars,
const struct radv_nir_compiler_options *options);