From 48952f041ba7bc2648bb5d94dcd270dc0f53bdff Mon Sep 17 00:00:00 2001 From: adjordjevic-TT Date: Fri, 10 Jan 2025 15:22:18 +0000 Subject: [PATCH 1/5] Save changes --- common/inc/cpack_common.h | 50 +++++++++++++++++++++++++++++++++++++ common/inc/cunpack_common.h | 35 ++++++++++++++++++++++++++ 2 files changed, 85 insertions(+) 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; + } + } From 0558d79f7850a21927e6ea937327284ac460f64e Mon Sep 17 00:00:00 2001 From: adjordjevic-TT Date: Fri, 17 Jan 2025 09:38:00 +0000 Subject: [PATCH 2/5] Changed read_pck_edge_offset to read_pack_edge_offset --- common/inc/cpack_common.h | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/common/inc/cpack_common.h b/common/inc/cpack_common.h index 519c9a5..8ee1d39 100644 --- a/common/inc/cpack_common.h +++ b/common/inc/cpack_common.h @@ -554,7 +554,7 @@ namespace ckernel::packer return dest.f; } - inline pck_edge_offset_t read_pck_edge_offset(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { + inline pck_edge_offset_t read_pack_edge_offset(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { pck_edge_offset_u edge = {.val=0}; From 0359cac5b80153440b6529981473a76fa73afb3c Mon Sep 17 00:00:00 2001 From: adjordjevic-TT Date: Fri, 17 Jan 2025 14:41:07 +0000 Subject: [PATCH 3/5] Changed read_pack_config to return array of config structures --- common/inc/cpack_common.h | 18 +++++++++++++++++- 1 file changed, 17 insertions(+), 1 deletion(-) diff --git a/common/inc/cpack_common.h b/common/inc/cpack_common.h index 8ee1d39..0b084cc 100644 --- a/common/inc/cpack_common.h +++ b/common/inc/cpack_common.h @@ -10,6 +10,8 @@ #include "ckernel_globals.h" #include "llk_defs.h" +#include + namespace ckernel::packer { constexpr uint replay_buf_offset = 16; // split replay buffer usage between fpu/sfpu @@ -521,7 +523,7 @@ namespace ckernel::packer // READERS FOR CONFIG STRUCTS - inline pack_config_t read_pack_config(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { + inline pack_config_t read_pack_config_helper(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { pack_config_u config = {.val = 0}; @@ -532,6 +534,20 @@ namespace ckernel::packer return config.f; } + inline std::array read_pack_config() { + std::array config_vec; + + // Get pointer to registers for current state ID + volatile uint tt_reg_ptr* cfg = get_cfg_pointer(); + + config_vec[0] = read_pack_config_helper(THCON_SEC0_REG1_Row_start_section_size_ADDR32, cfg); + config_vec[1] = read_pack_config_helper(THCON_SEC0_REG8_Row_start_section_size_ADDR32, cfg); + config_vec[2] = read_pack_config_helper(THCON_SEC1_REG1_Row_start_section_size_ADDR32, cfg); + config_vec[3] = read_pack_config_helper(THCON_SEC1_REG8_Row_start_section_size_ADDR32, cfg); + + return config_vec; + } + inline relu_config_t read_relu_config() { relu_config_u config; From 8f5ea78bc7e4b47ac4d0a89f31aaf54b243e2fdb Mon Sep 17 00:00:00 2001 From: adjordjevic-TT Date: Fri, 17 Jan 2025 15:10:42 +0000 Subject: [PATCH 4/5] Changed read_pack_counters to return array of config structs --- common/inc/cpack_common.h | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/common/inc/cpack_common.h b/common/inc/cpack_common.h index 0b084cc..cd73964 100644 --- a/common/inc/cpack_common.h +++ b/common/inc/cpack_common.h @@ -579,10 +579,24 @@ namespace ckernel::packer return edge.f; } - inline pack_counters_t read_pack_counters(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { + inline pack_counters_t read_pack_counters_helper(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; } + + inline std::array read_pack_counters() { + std::array config_vec; + + // Get pointer to registers for current state ID + volatile uint tt_reg_ptr* cfg = get_cfg_pointer(); + + config_vec[0] = read_pack_counters_helper(PACK_COUNTERS_SEC0_pack_per_xy_plane_ADDR32, cfg); + config_vec[1] = read_pack_counters_helper(PACK_COUNTERS_SEC1_pack_per_xy_plane_ADDR32, cfg); + config_vec[2] = read_pack_counters_helper(PACK_COUNTERS_SEC2_pack_per_xy_plane_ADDR32, cfg); + config_vec[3] = read_pack_counters_helper(PACK_COUNTERS_SEC3_pack_per_xy_plane_ADDR32, cfg); + + return config_vec; + } } From 579e7929f59da0a432674a1bcd55ea5b62898bef Mon Sep 17 00:00:00 2001 From: adjordjevic-TT Date: Fri, 17 Jan 2025 15:34:38 +0000 Subject: [PATCH 5/5] Changed read_pack_edge_offset to return array of config structs --- common/inc/cpack_common.h | 16 +++++++++++++++- 1 file changed, 15 insertions(+), 1 deletion(-) diff --git a/common/inc/cpack_common.h b/common/inc/cpack_common.h index cd73964..daa5c54 100644 --- a/common/inc/cpack_common.h +++ b/common/inc/cpack_common.h @@ -570,7 +570,7 @@ namespace ckernel::packer return dest.f; } - inline pck_edge_offset_t read_pack_edge_offset(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { + inline pck_edge_offset_t read_pack_edge_offset_helper(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { pck_edge_offset_u edge = {.val=0}; @@ -579,6 +579,20 @@ namespace ckernel::packer return edge.f; } + inline std::array read_pack_edge_offset() { + std::array edge_vec; + + // Get pointer to registers for current state ID + volatile uint tt_reg_ptr* cfg = get_cfg_pointer(); + + edge_vec[0] = read_pack_edge_offset_helper(PCK_EDGE_OFFSET_SEC0_mask_ADDR32, cfg); + edge_vec[1] = read_pack_edge_offset_helper(PCK_EDGE_OFFSET_SEC1_mask_ADDR32, cfg); + edge_vec[2] = read_pack_edge_offset_helper(PCK_EDGE_OFFSET_SEC2_mask_ADDR32, cfg); + edge_vec[3] = read_pack_edge_offset_helper(PCK_EDGE_OFFSET_SEC3_mask_ADDR32, cfg); + + return edge_vec; + } + inline pack_counters_t read_pack_counters_helper(uint32_t reg_addr, const volatile uint tt_reg_ptr* cfg) { pack_counters_u counters = {.val=0}; counters.val = cfg[reg_addr];