From 9c46e6fca323390f3cb74d6e865d2883a4fbd453 Mon Sep 17 00:00:00 2001 From: Samuel Pitoiset Date: Tue, 18 Aug 2020 18:39:20 +0200 Subject: [PATCH] aco: add a helper for building a trap handler shader MIME-Version: 1.0 Content-Type: text/plain; charset=utf8 Content-Transfer-Encoding: 8bit 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 Reviewed-by: Daniel Schürmann Part-of: --- .../compiler/aco_instruction_selection.cpp | 60 +++++++++++++++++++ src/amd/compiler/aco_ir.h | 3 + 2 files changed, 63 insertions(+) diff --git a/src/amd/compiler/aco_instruction_selection.cpp b/src/amd/compiler/aco_instruction_selection.cpp index 42b899e4c70..6f1f8b4e07e 100644 --- a/src/amd/compiler/aco_instruction_selection.cpp +++ b/src/amd/compiler/aco_instruction_selection.cpp @@ -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); +} } diff --git a/src/amd/compiler/aco_ir.h b/src/amd/compiler/aco_ir.h index 00a2e2596a3..18bc9bdd844 100644 --- a/src/amd/compiler/aco_ir.h +++ b/src/amd/compiler/aco_ir.h @@ -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); -- 2.30.2