Save changes

tenstorrent · Jan 10, 2025 · 48952f0 · 48952f0
1 parent 48e7139
commit 48952f0
Show file tree

Hide file tree

Showing 2 changed files with 85 additions and 0 deletions.
diff --git 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
@@ -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;
+   }
+
 }