Create Infrastructure to exactly calculate L1 Memory Usage for Conv2D #15088 #15455
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
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
2 times, most recently
from
2b3a64d to
d153dac
Compare
![github-actions[bot]](https://avatars.githubusercontent.com/in/15368?s=60&v=4){kind=small}
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
from
d153dac to
be2a301
Compare
sankarmanoj-tt
changed the title
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
from
be2a301 to
8d9ba30
Compare
sankarmanoj-tt
requested review from
mywoodstock,
shwetankTT,
abhullar-tt,
pgkeller,
aliuTT,
tt-aho,
tt-dma,
tt-asaigal and
ubcheema
as code owners
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
3 times, most recently
from
277ae36 to
73a0fe8
Compare
sankarmanoj-tt
requested review from
ayerofieiev-tt,
dmakoviichuk-tt,
cfjchu and
TT-BrianLiu
as code owners
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
from
d7e027d to
41c2469
Compare
abhullar-tt
approved these changes
Dec 18, 2024
ttnn/cpp/ttnn/operations/conv/conv2d/device/kernels/weights_reader_width_sharded.cpp
Outdated
Show resolved
Hide resolved
mywoodstock
reviewed
Jan 3, 2025
| uint32_t input_channels_alignment = 32; | ||
| bool deallocate_activation = false; | ||
| bool reallocate_halo_output = false; | ||
| uint32_t act_block_h_override = 0; // This argument is ignored when shard_layout == WIDTH_SHARDED. |
There was a problem hiding this comment.
Since this config is user facing, i think would be nicer to organize comments for each of the parameter -- something like:
/* this arg does that
*/
dtype arg = default_val;
There was a problem hiding this comment.
Fixed it. Is this okay?
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
from
bbfb12e to
9d1b535
Compare
mywoodstock
reviewed
Jan 8, 2025
| } | ||
|
|
||
| conv_op_l1_usage conv2d::estimate_L1_usage( |
There was a problem hiding this comment.
This is exact calculation, right? Perhaps rename to calculate_L1_usage?
There was a problem hiding this comment.
Yes. I have renamed it.
| device->get_memory_allocation_statistics(tt::tt_metal::BufferType::L1).total_allocated_bytes; | ||
| auto actual_cb_size = program_with_cbs.program.get_cb_memory_size(); | ||
|
|
||
| auto [calc_output_size, calc_CB_size] = estimate_L1_usage( |
There was a problem hiding this comment.
Will this have e2e perf degradation?
There was a problem hiding this comment.
The function itself just does a handful of calculations, and so should take very little time. It's invoked when the program object is created, in operation::launch_op. AFAIK, that should only be called the first time, and subsequent calls should hit the cache.
mywoodstock
approved these changes
Jan 8, 2025
There was a problem hiding this comment.
@sankarmanoj-tt please wait merging until @pavlejosipovic also approves.
ttnn/cpp/ttnn/operations/conv/conv2d/prepare_conv2d_weights.hpp
Outdated
Show resolved
Hide resolved
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
from
42751fc to
4435b19
Compare
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
from
7c7ac78 to
cd60c9f
Compare
pavlejosipovic
approved these changes
Jan 9, 2025
ayerofieiev-tt
approved these changes
Jan 11, 2025
sankarmanoj-tt
force-pushed
the
smanoj/estimate_l1_conv
branch
from
a01be9d to
93461c8
Compare
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
#15088
Problem description
Accurate ahead of time calculation of L1 usage is necessary for enabling lots of features and optimizations.
What's changed
Implemented code to calculate L1 usage, and cross-verify it with the actual L1 usage. The L1 usage of an op consists of the output tensor and the circular buffers.
The L1 usage for the output tensor can be obtained by checking the global L1 memory allocation stats before and after the op.
The L1 usage by Circular Buffers is calculated from the program object. A new method was added to the program object to be able to get the total memory usage by Circular Buffer before memory allocation.
Checklist