Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Added config registers readers #63

Open
wants to merge 2 commits into
base: main
Choose a base branch
from
Open

Added config registers readers #63

Show file tree
Hide file tree
Changes from 1 commit
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
48952f0
Save changes
adjordjevic-TT Jan 10, 2025
0558d79
Changed read_pck_edge_offset to read_pack_edge_offset
adjordjevic-TT Jan 17, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
50 changes: 50 additions & 0 deletions common/inc/cpack_common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

For other functions this is not passed as input since it is known at compile time.
Can you use the same thing here?


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();
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why not keep the same approach as for read_pack_config? Where cfg is passed as argument?

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) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Let's please not use abbreviation. If anything, we should remove the existing ones. pck -> pack


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) {
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Same comment about arguments as above. cfg could be obtained inside function and reg_addr should be known at compile time. If true, let's use that.
It also reduces possibility for user errors, as currently user can pass a wrong reg_addr and read pack_config instead of counters, for example.

pack_counters_u counters = {.val=0};
counters.val = cfg[reg_addr];

return counters.f;
}
}
35 changes: 35 additions & 0 deletions common/inc/cunpack_common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -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];
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Will this correctly populate just the FPU rounding bit in the config? Or the whole config, and just the register name is unfortunate?

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

This should populate whole config. ALU_ROUNDING_MODE_Fpu_srnd_en_ADDR32 is starting address of the whole config (since fpu_srnd_en is first field in structure). As I understood, we pass to cfg starting address and then it returns 32 bits starting from there.


return config.f;
}

}
Loading