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

#11830: Move ALLOCATOR_ALIGNMENT into the allocator #13459

Merged
merged 1 commit into from
Oct 7, 2024
Merged

#11830: Move ALLOCATOR_ALIGNMENT into the allocator #13459

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
2 changes: 1 addition & 1 deletion tech_reports/prog_examples/shard_data_rm/shard_data_rm.md
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -63,7 +63,7 @@ uint32_t shard_size = shard_height * shard_width;
uint32_t input_unit_size = sizeof(uint32_t);
uint32_t shard_width_bytes = shard_width * data_size;
uint32_t num_units_per_row = shard_width * input_unit_size;
uint32_t padded_offset_bytes = align(input_unit_size, ALLOCATOR_ALIGNMENT);
uint32_t padded_offset_bytes = align(input_unit_size, device->get_allocator_alignment());
```

In order to shard the correct data segments to the respective core, we indicate the shard height, width, size, and other data for the kernel function.
Expand Down
4 changes: 2 additions & 2 deletions tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -591,9 +591,9 @@ inline void generate_random_paged_payload(Device *device,
log_debug(tt::LogTest, "Starting {} w/ is_dram: {} start_page: {} words_per_page: {}", __FUNCTION__, is_dram, start_page, words_per_page);

// Note: the dst address marches in unison regardless of whether or not a core is written to
uint32_t page_size_alignment_bytes = device->get_allocator_alignment();
for (uint32_t page_id = start_page; page_id < start_page + cmd.write_paged.pages; page_id++) {

constexpr uint32_t page_size_alignment_bytes = ALLOCATOR_ALIGNMENT;
CoreCoord bank_core;
uint32_t bank_id = page_id % num_banks;
uint32_t bank_offset = align(cmd.write_paged.page_size, page_size_alignment_bytes) * (page_id / num_banks);
Expand Down Expand Up @@ -875,7 +875,7 @@ inline void gen_dispatcher_paged_write_cmd(Device *device,
uint32_t page_size,
uint32_t pages) {

constexpr uint32_t page_size_alignment_bytes = ALLOCATOR_ALIGNMENT;
uint32_t page_size_alignment_bytes = device->get_allocator_alignment();
uint32_t num_banks = device->num_banks(is_dram ? BufferType::DRAM : BufferType::L1);
CoreType core_type = is_dram ? CoreType::DRAM : CoreType::WORKER;

Expand Down
33 changes: 18 additions & 15 deletions tests/tt_metal/tt_metal/unit_tests/allocator/test_l1_banking_allocator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -10,20 +10,30 @@
#include "tt_metal/detail/tt_metal.hpp"
#include "tt_metal/host_api.hpp"

// TODO: Uplift to DeviceFixture once it does not skip GS
TEST_F(BasicFixture, TestL1BuffersAllocatedTopDown) {
tt::tt_metal::Device *device = tt::tt_metal::CreateDevice(0, 1, 0);

std::vector<uint32_t> alloc_sizes = {32 * 1024, 64 * 1024, 128 * 1024};
size_t total_size_bytes = 0;
namespace unit_tests::test_l1_banking_allocator {

uint64_t get_alloc_limit(const tt::tt_metal::Device *device) {
const metal_SocDescriptor &soc_desc = tt::Cluster::instance().get_soc_desc(device->id());
uint32_t l1_unreserved_base = device->get_base_allocator_addr(tt::tt_metal::HalMemType::L1);
CoreType dispatch_core_type = dispatch_core_manager::instance().get_dispatch_core_type(device->id());
auto storage_core_bank_size = tt::get_storage_core_bank_size(device->id(), device->num_hw_cqs(), dispatch_core_type);
const uint32_t allocator_alignment = device->get_allocator_alignment();
const uint32_t interleaved_l1_bank_size = storage_core_bank_size.has_value() ? storage_core_bank_size.value() : (soc_desc.worker_l1_size - l1_unreserved_base);
uint32_t storage_core_unreserved_base = ((MEM_MAILBOX_BASE + ALLOCATOR_ALIGNMENT - 1) / ALLOCATOR_ALIGNMENT) * ALLOCATOR_ALIGNMENT;
uint32_t storage_core_unreserved_base = ((MEM_MAILBOX_BASE + allocator_alignment - 1) / allocator_alignment) * allocator_alignment;
uint64_t alloc_limit = interleaved_l1_bank_size - storage_core_unreserved_base;
return alloc_limit;
}

} // namespace unit_tests::test_l1_banking_allocator

// TODO: Uplift to DeviceFixture once it does not skip GS
TEST_F(BasicFixture, TestL1BuffersAllocatedTopDown) {
tt::tt_metal::Device *device = tt::tt_metal::CreateDevice(0, 1, 0);

std::vector<uint32_t> alloc_sizes = {32 * 1024, 64 * 1024, 128 * 1024};
size_t total_size_bytes = 0;

uint64_t alloc_limit = unit_tests::test_l1_banking_allocator::get_alloc_limit(device);

std::vector<std::unique_ptr<Buffer>> buffers;
int alloc_size_idx = 0;
Expand All @@ -47,14 +57,7 @@ TEST_F(BasicFixture, TestL1BuffersAllocatedTopDown) {
// TODO: Uplift to DeviceFixture once it does not skip GS
TEST_F(BasicFixture, TestL1BuffersDoNotGrowBeyondBankSize) {
tt::tt_metal::Device *device = tt::tt_metal::CreateDevice(0, 1, 0);

const metal_SocDescriptor &soc_desc = tt::Cluster::instance().get_soc_desc(device->id());
uint32_t l1_unreserved_base = device->get_base_allocator_addr(tt::tt_metal::HalMemType::L1);
CoreType dispatch_core_type = dispatch_core_manager::instance().get_dispatch_core_type(device->id());
auto storage_core_bank_size = tt::get_storage_core_bank_size(device->id(), device->num_hw_cqs(), dispatch_core_type);
const uint32_t interleaved_l1_bank_size = storage_core_bank_size.has_value() ? storage_core_bank_size.value() : (soc_desc.worker_l1_size - l1_unreserved_base);
uint32_t storage_core_unreserved_base = ((MEM_MAILBOX_BASE + ALLOCATOR_ALIGNMENT - 1) / ALLOCATOR_ALIGNMENT) * ALLOCATOR_ALIGNMENT;
uint64_t alloc_limit = interleaved_l1_bank_size - storage_core_unreserved_base;
uint64_t alloc_limit = unit_tests::test_l1_banking_allocator::get_alloc_limit(device);

tt::tt_metal::InterleavedBufferConfig l1_config{
.device=device,
Expand Down
6 changes: 0 additions & 6 deletions tt_metal/hostdevcommon/common_runtime_address_map.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -19,12 +19,6 @@
constexpr static std::uint32_t DRAM_BARRIER_BASE = 0;
constexpr static std::uint32_t DRAM_BARRIER_SIZE = ((sizeof(uint32_t) + DRAM_ALIGNMENT - 1) / DRAM_ALIGNMENT) * DRAM_ALIGNMENT;

// Take max alignment to satisfy NoC rd/wr constraints
// Tensix/Eth -> PCIe/DRAM src and dst addrs must be L1_ALIGNMENT aligned
// PCIe/DRAM -> Tensix/Eth src and dst addrs must be DRAM_ALIGNMENT aligned
// Tensix/Eth <-> Tensix/Eth src and dst addrs must be L1_ALIGNMENT aligned
constexpr static std::uint32_t ALLOCATOR_ALIGNMENT = DRAM_ALIGNMENT >= L1_ALIGNMENT ? DRAM_ALIGNMENT : L1_ALIGNMENT;

// TODO: move these out of the memory map into profiler code
constexpr static std::uint32_t PROFILER_OP_SUPPORT_COUNT = 1000;
constexpr static std::uint32_t PROFILER_FULL_HOST_VECTOR_SIZE_PER_RISC = kernel_profiler::PROFILER_L1_MARKER_UINT32_SIZE * (kernel_profiler::PROFILER_L1_PROGRAM_ID_COUNT + kernel_profiler::PROFILER_L1_GUARANTEED_MARKER_COUNT + kernel_profiler::PROFILER_L1_OP_MIN_OPTIONAL_MARKER_COUNT) * PROFILER_OP_SUPPORT_COUNT;
Expand Down
6 changes: 3 additions & 3 deletions tt_metal/impl/allocator/allocator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -215,7 +215,7 @@ void init_one_bank_per_channel(Allocator &allocator, const AllocatorConfig &allo
bank_offsets.at(channel_id) = static_cast<int32_t>(alloc_config.dram_bank_offsets.at(channel_id));
}
allocator.dram_manager =
BankManager(BufferType::DRAM, bank_offsets, dram_bank_size, ALLOCATOR_ALIGNMENT, alloc_config.dram_unreserved_base);
BankManager(BufferType::DRAM, bank_offsets, dram_bank_size, alloc_config.alignment, alloc_config.dram_unreserved_base);
for (uint32_t bank_id = 0; bank_id < alloc_config.num_dram_channels; bank_id++) {
CoreCoord logical_core = CoreCoord{bank_id, 0};
allocator.bank_id_to_dram_channel.insert({bank_id, bank_id});
Expand All @@ -228,7 +228,7 @@ void init_one_bank_per_channel(Allocator &allocator, const AllocatorConfig &allo
BufferType::TRACE,
bank_offsets,
alloc_config.trace_region_size,
ALLOCATOR_ALIGNMENT,
alloc_config.alignment,
dram_bank_size + alloc_config.dram_unreserved_base);
for (uint32_t bank_id = 0; bank_id < alloc_config.num_dram_channels; bank_id++) {
CoreCoord logical_core = CoreCoord{bank_id, 0};
Expand All @@ -244,7 +244,7 @@ void init_one_bank_per_l1(Allocator &allocator, const AllocatorConfig &alloc_con
// Space up to L1 unreserved base is reserved for risc binaries, kernel args, debug and perf monitoring tools
DeviceAddr l1_bank_size = alloc_config.worker_l1_size - alloc_config.l1_unreserved_base;
std::vector<int64_t> bank_offsets(num_l1_banks, 0);
allocator.l1_manager = BankManager(BufferType::L1, bank_offsets, l1_bank_size, ALLOCATOR_ALIGNMENT, alloc_config.l1_unreserved_base);
allocator.l1_manager = BankManager(BufferType::L1, bank_offsets, l1_bank_size, alloc_config.alignment, alloc_config.l1_unreserved_base);

uint32_t bank_id = 0;
for (uint32_t y = 0; y < alloc_config.worker_grid_size.y; y++) {
Expand Down
1 change: 1 addition & 0 deletions tt_metal/impl/allocator/allocator_types.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,6 +50,7 @@ struct AllocatorConfig {
std::unordered_map<int, int> worker_log_to_physical_routing_y = {};
BankMapping l1_bank_remap = {}; // for remapping which l1 bank points to which bank if we assume normal row-major assignment
CoreCoord compute_grid_size = {};
uint32_t alignment = 0;
void reset();
~AllocatorConfig() { reset(); }
};
Expand Down
6 changes: 3 additions & 3 deletions tt_metal/impl/allocator/l1_banking_allocator.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -171,7 +171,7 @@ void init_compute_and_storage_l1_bank_manager(Allocator &allocator, const Alloca
num_banks.total);

// Storage only cores only need to reserve mailbox space to hold barriers
constexpr uint32_t storage_core_unreserved_base = ((MEM_MAILBOX_BASE + ALLOCATOR_ALIGNMENT - 1) / ALLOCATOR_ALIGNMENT) * ALLOCATOR_ALIGNMENT;
uint32_t storage_core_unreserved_base = ((MEM_MAILBOX_BASE + alloc_config.alignment - 1) / alloc_config.alignment) * alloc_config.alignment;

// There is only l1_bank_size bytes available for L1 buffers to be allocated in
uint64_t l1_bank_size = alloc_config.storage_core_bank_size.has_value()
Expand All @@ -181,7 +181,7 @@ void init_compute_and_storage_l1_bank_manager(Allocator &allocator, const Alloca
uint64_t allocatable_l1_size =
static_cast<uint64_t>(alloc_config.worker_l1_size) - alloc_config.l1_unreserved_base - alloc_config.l1_small_size;
// Assuming top down allocation for L1 buffers so the allocatable memory space is the top l1_bank_size bytes of L1
allocator.l1_manager = BankManager(BufferType::L1, bank_id_to_bank_offset, allocatable_l1_size, interleaved_address_limit, ALLOCATOR_ALIGNMENT, alloc_config.l1_unreserved_base);
allocator.l1_manager = BankManager(BufferType::L1, bank_id_to_bank_offset, allocatable_l1_size, interleaved_address_limit, alloc_config.alignment, alloc_config.l1_unreserved_base);

uint64_t small_interleaved_address_limit = alloc_config.worker_l1_size - alloc_config.l1_small_size;
uint64_t small_alloc_offset = alloc_config.l1_unreserved_base + allocatable_l1_size;
Expand All @@ -193,7 +193,7 @@ void init_compute_and_storage_l1_bank_manager(Allocator &allocator, const Alloca
small_bank_id_to_bank_offset,
alloc_config.l1_small_size,
small_interleaved_address_limit,
ALLOCATOR_ALIGNMENT,
alloc_config.alignment,
small_alloc_offset);
}

Expand Down
2 changes: 2 additions & 0 deletions tt_metal/impl/buffers/buffer.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -313,6 +313,8 @@ DeviceAddr Buffer::page_address(uint32_t bank_id, uint32_t page_index) const {
return translate_page_address(offset, bank_id);
}

uint32_t Buffer::alignment() const { return this->device_->get_allocator_alignment(); }

DeviceAddr Buffer::sharded_page_address(uint32_t bank_id, uint32_t page_index) const {
TT_ASSERT(is_sharded(this->buffer_layout()));
int pages_offset_within_bank = page_index % shard_spec().size();
Expand Down
2 changes: 1 addition & 1 deletion tt_metal/impl/buffers/buffer.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -225,7 +225,7 @@ class Buffer {

DeviceAddr page_address(uint32_t bank_id, uint32_t page_index) const;

uint32_t alignment() const { return ALLOCATOR_ALIGNMENT; }
uint32_t alignment() const;

DeviceAddr aligned_page_size() const { return align(page_size_, this->alignment());}

Expand Down
21 changes: 16 additions & 5 deletions tt_metal/impl/device/device.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -209,6 +209,10 @@ void Device::initialize_allocator(size_t l1_small_size, size_t trace_region_size
const metal_SocDescriptor &soc_desc = tt::Cluster::instance().get_soc_desc(this->id_);
CoreType dispatch_core_type = dispatch_core_manager::instance().get_dispatch_core_type(this->id_);
// Construct allocator config from soc_desc
// Take max alignment to satisfy NoC rd/wr constraints
// Tensix/Eth -> PCIe/DRAM src and dst addrs must be L1_ALIGNMENT aligned
// PCIe/DRAM -> Tensix/Eth src and dst addrs must be DRAM_ALIGNMENT aligned
// Tensix/Eth <-> Tensix/Eth src and dst addrs must be L1_ALIGNMENT aligned
AllocatorConfig config(
{.num_dram_channels = static_cast<size_t>(soc_desc.get_num_dram_channels()),
.dram_bank_size = soc_desc.dram_bank_size,
Expand All @@ -224,13 +228,14 @@ void Device::initialize_allocator(size_t l1_small_size, size_t trace_region_size
.worker_log_to_physical_routing_x = soc_desc.worker_log_to_physical_routing_x,
.worker_log_to_physical_routing_y = soc_desc.worker_log_to_physical_routing_y,
.l1_bank_remap = l1_bank_remap,
.compute_grid_size = this->compute_with_storage_grid_size()});
.compute_grid_size = this->compute_with_storage_grid_size(),
.alignment = std::max(hal.get_alignment(HalMemType::DRAM), hal.get_alignment(HalMemType::L1))});
TT_FATAL(config.l1_small_size < (config.storage_core_bank_size.has_value() ? config.storage_core_bank_size.value() : config.worker_l1_size - config.l1_unreserved_base),
"Reserved size must be less than bank size");
TT_FATAL(
config.l1_small_size % ALLOCATOR_ALIGNMENT == 0,
"Reserved size must be aligned to ALLOCATOR_ALIGNMENT {}",
ALLOCATOR_ALIGNMENT);
config.l1_small_size % config.alignment == 0,
"Reserved size must be aligned to allocator alignment {}",
config.alignment);
// Initialize dram_offsets from soc_descriptor
for (auto channel = 0; channel < soc_desc.get_num_dram_channels(); channel++) {
config.dram_bank_offsets.push_back(soc_desc.get_address_offset(channel));
Expand Down Expand Up @@ -3100,6 +3105,11 @@ allocator::Statistics Device::get_memory_allocation_statistics(const BufferType
return allocator::get_statistics(*this->allocator_, buffer_type);
}

uint32_t Device::get_allocator_alignment() const {
this->check_allocator_is_initialized();
return this->allocator_->config.alignment;
}

size_t Device::get_l1_small_size() const {
this->check_allocator_is_initialized();
return this->allocator_->config.l1_small_size;
Expand Down Expand Up @@ -3327,7 +3337,8 @@ void Device::generate_device_headers(const std::string &path) const
dram_noc_coord_per_bank,
dram_offsets_per_bank,
l1_noc_coord_per_bank,
l1_offset_per_bank
l1_offset_per_bank,
this->allocator_->config.alignment
);
}

Expand Down
2 changes: 2 additions & 0 deletions tt_metal/impl/device/device.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -169,6 +169,8 @@ class Device {

allocator::Statistics get_memory_allocation_statistics(const BufferType &buffer_type) const;

uint32_t get_allocator_alignment() const;

size_t get_l1_small_size() const;

void dump_memory_blocks(const BufferType &buffer_type, std::ofstream &out) const;
Expand Down
12 changes: 8 additions & 4 deletions tt_metal/jit_build/genfiles.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -524,7 +524,8 @@ std::string generate_bank_to_noc_coord_descriptor_string(
std::vector<CoreCoord>& dram_bank_map,
std::vector<int32_t>& dram_bank_offset_map,
std::vector<CoreCoord>& l1_bank_map,
std::vector<int32_t>& l1_bank_offset_map) {
std::vector<int32_t>& l1_bank_offset_map,
uint32_t allocator_alignment) {
stringstream ss;
bool is_dram_pow2 = ceil(log2(dram_bank_map.size())) == log2(dram_bank_map.size());
bool is_l1_pow2 = ceil(log2(l1_bank_map.size())) == log2(l1_bank_map.size());
Expand All @@ -546,7 +547,8 @@ std::string generate_bank_to_noc_coord_descriptor_string(
ss << "#include <noc/noc_parameters.h>" << endl;
ss << endl;

ss << "#define LOG_BASE_2_OF_ALLOCATOR_ALIGNMENT " << std::bit_width(ALLOCATOR_ALIGNMENT) - 1 << endl;
ss << "#define ALLOCATOR_ALIGNMENT " << allocator_alignment << endl;
ss << "#define LOG_BASE_2_OF_ALLOCATOR_ALIGNMENT " << std::bit_width(allocator_alignment) - 1 << endl;
ss << "#define NUM_DRAM_BANKS " << dram_bank_map.size() << endl;
ss << "#define NUM_L1_BANKS " << l1_bank_map.size() << endl;

Expand Down Expand Up @@ -634,13 +636,15 @@ void jit_build_genfiles_bank_to_noc_coord_descriptor(
std::vector<CoreCoord>& dram_bank_map,
std::vector<int32_t>& dram_bank_offset_map,
std::vector<CoreCoord>& l1_bank_map,
std::vector<int32_t>& l1_bank_offset_map) {
std::vector<int32_t>& l1_bank_offset_map,
uint32_t allocator_alignment) {
string output_string = generate_bank_to_noc_coord_descriptor_string(
grid_size,
dram_bank_map,
dram_bank_offset_map,
l1_bank_map,
l1_bank_offset_map);
l1_bank_offset_map,
allocator_alignment);

fs::create_directories(path + "/brisc");
ofstream file_stream_br(path + "/brisc/generated_bank_to_noc_coord_mapping.h");
Expand Down
3 changes: 2 additions & 1 deletion tt_metal/jit_build/genfiles.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -27,7 +27,8 @@ void jit_build_genfiles_bank_to_noc_coord_descriptor(
std::vector<CoreCoord>& dram_bank_map,
std::vector<int32_t>& dram_bank_offset_map,
std::vector<CoreCoord>& l1_bank_map,
std::vector<int32_t>& l1_bank_offset_map);
std::vector<int32_t>& l1_bank_offset_map,
uint32_t allocator_alignment);

void jit_build_genfiles_descriptors(const JitBuildEnv& env, JitBuildOptions& options);

Expand Down
3 changes: 2 additions & 1 deletion tt_metal/llrt/blackhole/bh_hal_idle_eth.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -23,6 +23,7 @@ namespace tt_metal {
HalCoreInfoType create_idle_eth_mem_map() {

constexpr uint32_t num_proc_per_idle_eth_core = 1;
uint32_t max_alignment = std::max(DRAM_ALIGNMENT, L1_ALIGNMENT);

std::vector<DeviceAddr> mem_map_bases;

Expand All @@ -33,7 +34,7 @@ HalCoreInfoType create_idle_eth_mem_map() {
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::DPRINT)] = GET_IERISC_MAILBOX_ADDRESS_HOST(dprint_buf);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::PROFILER)] = GET_IERISC_MAILBOX_ADDRESS_HOST(profiler);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::KERNEL_CONFIG)] = IDLE_ERISC_L1_KERNEL_CONFIG_BASE;
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::UNRESERVED)] = ((L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE - 1) | (ALLOCATOR_ALIGNMENT - 1)) + 1; // TODO: this is wrong, need idle eth specific value
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::UNRESERVED)] = ((L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE - 1) | (max_alignment - 1)) + 1; // TODO: this is wrong, need idle eth specific value
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::CORE_INFO)] = GET_IERISC_MAILBOX_ADDRESS_HOST(core_info);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::GO_MSG)] = GET_IERISC_MAILBOX_ADDRESS_HOST(go_message);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::LAUNCH_MSG_BUFFER_RD_PTR)] = GET_IERISC_MAILBOX_ADDRESS_HOST(launch_msg_rd_ptr);
Expand Down
3 changes: 2 additions & 1 deletion tt_metal/llrt/blackhole/bh_hal_tensix.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -21,6 +21,7 @@ namespace tt_metal {
HalCoreInfoType create_tensix_mem_map() {

constexpr uint32_t num_proc_per_tensix_core = 5;
uint32_t max_alignment = std::max(DRAM_ALIGNMENT, L1_ALIGNMENT);

std::vector<DeviceAddr> mem_map_bases;

Expand All @@ -31,7 +32,7 @@ HalCoreInfoType create_tensix_mem_map() {
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::DPRINT)] = GET_MAILBOX_ADDRESS_HOST(dprint_buf);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::PROFILER)] = GET_MAILBOX_ADDRESS_HOST(profiler);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::KERNEL_CONFIG)] = L1_KERNEL_CONFIG_BASE;
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::UNRESERVED)] = ((L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE - 1) | (ALLOCATOR_ALIGNMENT - 1)) + 1;
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::UNRESERVED)] = ((L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE - 1) | (max_alignment - 1)) + 1;
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::CORE_INFO)] = GET_MAILBOX_ADDRESS_HOST(core_info);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::GO_MSG)] = GET_MAILBOX_ADDRESS_HOST(go_message);
mem_map_bases[utils::underlying_type<HalMemAddrType>(HalMemAddrType::LAUNCH_MSG_BUFFER_RD_PTR)] = GET_MAILBOX_ADDRESS_HOST(launch_msg_rd_ptr);
Expand Down
Loading
Loading