diff --git a/common/inc/cpack_common.h b/common/inc/cpack_common.h index 38deacf..519c9a5 100644 --- a/common/inc/cpack_common.h +++ b/common/inc/cpack_common.h @@ -519,4 +519,54 @@ namespace ckernel::packer TTI_STOREIND (1, 0, p_ind::LD_16B, LO_16(0), p_ind::INC_NONE, p_gpr_pack::TILE_HEADER, p_gpr_pack::OUTPUT_ADDR); } + // READERS FOR CONFIG STRUCTS + + inline pack_config_t read_pack_config(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { + + pack_config_u config = {.val = 0}; + + config.val[0] = cfg[reg_addr]; + config.val[1] = cfg[reg_addr + 1]; + config.val[2] = cfg[reg_addr + 2]; + + return config.f; + } + + inline relu_config_t read_relu_config() { + + relu_config_u config; + + // Get pointer to registers for current state ID + volatile uint tt_reg_ptr *cfg = get_cfg_pointer(); + config.val[0] = cfg[ALU_ACC_CTRL_Zero_Flag_disabled_src_ADDR32]; + + return config.r; + } + + inline dest_rd_ctrl_t read_dest_rd_ctrl() { + dest_rd_ctrl_u dest; + + // Get pointer to registers for current state ID + volatile uint tt_reg_ptr *cfg = get_cfg_pointer(); + + dest.val = cfg[PCK_DEST_RD_CTRL_Read_32b_data_ADDR32]; + + return dest.f; + } + + inline pck_edge_offset_t read_pck_edge_offset(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { + + pck_edge_offset_u edge = {.val=0}; + + edge.val = cfg[reg_addr]; + + return edge.f; + } + + inline pack_counters_t read_pack_counters(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { + pack_counters_u counters = {.val=0}; + counters.val = cfg[reg_addr]; + + return counters.f; + } } diff --git a/common/inc/cunpack_common.h b/common/inc/cunpack_common.h index 41ac495..d1e36ce 100644 --- a/common/inc/cunpack_common.h +++ b/common/inc/cunpack_common.h @@ -469,4 +469,39 @@ namespace ckernel::unpacker } + // READERS FOR STRUCTS + + inline unpack_tile_descriptor_t read_unpack_tile_descriptor() { + unpack_tile_descriptor_u tile_descriptor; + volatile uint tt_reg_ptr *cfg = get_cfg_pointer(); + + tile_descriptor.val[0] = cfg[THCON_SEC0_REG0_TileDescriptor_ADDR32]; + tile_descriptor.val[1] = cfg[THCON_SEC0_REG0_TileDescriptor_ADDR32 + 1]; + tile_descriptor.val[2] = cfg[THCON_SEC1_REG0_TileDescriptor_ADDR32]; + tile_descriptor.val[3] = cfg[THCON_SEC1_REG0_TileDescriptor_ADDR32 + 1]; + + return tile_descriptor.f; + } + + inline unpack_config_t read_unpack_config() { + unpack_config_u config; + volatile uint tt_reg_ptr *cfg = get_cfg_pointer(); + + config.val[0] = cfg[THCON_SEC0_REG2_Out_data_format_ADDR32]; + config.val[1] = cfg[THCON_SEC0_REG2_Out_data_format_ADDR32 + 1]; + config.val[2] = cfg[THCON_SEC1_REG2_Out_data_format_ADDR32]; + config.val[3] = cfg[THCON_SEC1_REG2_Out_data_format_ADDR32 + 1]; + + return config.f; + } + + inline alu_config_t read_alu_config() { + alu_config_u config; + volatile uint tt_reg_ptr *cfg = get_cfg_pointer(); + + config.val = cfg[ALU_ROUNDING_MODE_Fpu_srnd_en_ADDR32]; + + return config.f; + } + }