#13609: Uplift dram and l1 allocators to use dram/l1 specific alignment #13762
+117
−50
Conversation
This file contains bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
llongTT
requested review from
ayerofieiev-tt,
dmakoviichuk-tt,
rfurko-tt,
TT-BrianLiu,
razorback3,
dongjin-na,
bbradelTT,
ntarafdar,
sjameelTT,
jaykru-tt,
yugi957,
jvegaTT and
llongTT
as code owners
bbradelTT
approved these changes
Dec 20, 2024
| @@ -1653,7 +1653,7 @@ operation::ProgramWithCallbacks pad_rm_sharded_width_only( | |||
|
|
|||
| // FIXME: assumes that this was sharded using DRAM alignment so that gaps are left in the tensor. | |||
There was a problem hiding this comment.
Should this comment be udpated as well?
| @@ -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, device->get_allocator_alignment()); | |||
| uint32_t padded_offset_bytes = align(input_unit_size, device->get_allocator_alignment(BufferType::L1)); | |||
There was a problem hiding this comment.
It might be good to add more info about when you need or when the infra uses different alignments and why.
tt-aho
reviewed
Dec 20, 2024
| @@ -676,7 +676,7 @@ operation::ProgramWithCallbacks transpose_hc_multi_core( | |||
| // TODO: noc_async_write only require 16B alignment for both DRAM and L1 for Blackhole, so instead of reading in | |||
| // face-lines from C tiles to form a single tile, we can load a single tile and then write out its face-lines to C | |||
| // tiles | |||
| uint32_t alignment = dst_buffer->buffer_type() == tt::tt_metal::BufferType::DRAM ? DRAM_ALIGNMENT : L1_ALIGNMENT; | |||
| uint32_t alignment = device->get_allocator_alignment(dst_buffer->buffer_type()); | |||
There was a problem hiding this comment.
Use dst_buffer->alignment() instead
| @@ -1653,7 +1653,7 @@ operation::ProgramWithCallbacks pad_rm_sharded_width_only( | |||
|
|
|||
| // FIXME: assumes that this was sharded using DRAM alignment so that gaps are left in the tensor. | |||
| // if this changes, we should change the stick step to be 16B (L1 alignment). | |||
| auto dram_alignment_bytes = tt::tt_metal::hal.get_alignment(tt::tt_metal::HalMemType::DRAM); | |||
| auto dram_alignment_bytes = tt::tt_metal::hal.get_alignment(tt::tt_metal::HalMemType::L1); | |||
There was a problem hiding this comment.
Should be renamed to l1_alignment_bytes?
tt_metal/hw/inc/dataflow_api.h
Outdated
| @@ -921,9 +941,10 @@ struct InterleavedPow2AddrGen { | |||
| const uint32_t bank_base_address; | |||
| const uint32_t log_base_2_of_page_size; // WARNING: This struct is used for optimized get_noc_addr in which case | |||
| // you know that bank_unit_size is a power of 2 | |||
| const uint32_t aligned_log_base_2_of_page_size = this->log_base_2_of_page_size > LOG_BASE_2_OF_ALLOCATOR_ALIGNMENT | |||
| const uint32_t log_base_2_of_allocator_alignment = interleaved_addr_gen::get_log_base2_of_allocator_alignment<DRAM>(); | |||
This comment was marked as resolved.
This comment was marked as resolved.
Sorry, something went wrong.
tt_metal/hw/inc/dataflow_api.h
Outdated
| @@ -1019,9 +1040,10 @@ template <bool DRAM> | |||
| struct InterleavedPow2AddrGenFast { | |||
| uint32_t bank_base_address; // Base address for the whole tensor. | |||
| const uint32_t log_base_2_of_page_size; // Num bytes in bank unit. | |||
| const uint32_t aligned_log_base_2_of_page_size = this->log_base_2_of_page_size > LOG_BASE_2_OF_ALLOCATOR_ALIGNMENT | |||
| const uint32_t log_base_2_of_allocator_alignment = interleaved_addr_gen::get_log_base2_of_allocator_alignment<DRAM>(); | |||
This comment was marked as resolved.
This comment was marked as resolved.
Sorry, something went wrong.
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.
This suggestion is invalid because no changes were made to the code.
Suggestions cannot be applied while the pull request is closed.
Suggestions cannot be applied while viewing a subset of changes.
Only one suggestion per line can be applied in a batch.
Add this suggestion to a batch that can be applied as a single commit.
Applying suggestions on deleted lines is not supported.
You must change the existing code in this line in order to create a valid suggestion.
Outdated suggestions cannot be applied.
This suggestion has been applied or marked resolved.
Suggestions cannot be applied from pending reviews.
Suggestions cannot be applied on multi-line comments.
Suggestions cannot be applied while the pull request is queued to merge.
Suggestion cannot be applied right now. Please check back later.
Ticket
#13609
Problem description
Using the max of DRAM and L1 alignment for both DRAM and L1 buffers was causing pcc mismatches in i2s and s2i.
What's changed
Use L1/DRAM specific alignment for respective allocations. This will require some ops to be uplifted to handle re-alignment
@yugaoTT and @ntarafdar to add corresponding op changes
Checklist
Below post commits were triggered 12/03