Fix cb allocation errors for halo and conv2d #16190
Merged
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
pavlejosipovic
requested review from
esmalTT,
uaydonat,
mywoodstock,
shwetankTT,
tt-aho and
sankarmanoj-tt
as code owners
pavlejosipovic
commented
Dec 19, 2024
mywoodstock
approved these changes
Dec 20, 2024
| @@ -462,6 +462,7 @@ def __call__(self, x, move_input_tensor_to_device=True): | |||
| ttnn.deallocate(c4_residual) | |||
| x = self.upblock2(x, c3_residual) | |||
| ttnn.deallocate(c3_residual) | |||
| c2_residual = ttnn.to_memory_config(c2_residual, ttnn.L1_MEMORY_CONFIG) | |||
There was a problem hiding this comment.
Is this the reason for regression on unet perf?
There was a problem hiding this comment.
It would be good to figure out why this is needed
There was a problem hiding this comment.
yes it is, I don't have bandwidth to dig into why is this helping
ttnn/cpp/ttnn/operations/conv/conv2d/prepare_conv2d_weights.cpp
Outdated
Show resolved
Hide resolved
tt-aho
approved these changes
Dec 23, 2024
esmalTT
approved these changes
Dec 23, 2024
uaydonat
approved these changes
Dec 23, 2024
pavlejosipovic
force-pushed
the
pjosipovic/fix_cb_asserts
branch
2 times, most recently
from
c6b8fb8 to
a9f2805
Compare
Halo op has a problem where manipulating tt::tt_metal::LegacyShape in compute_output_specs is not doing the expected thing (value is unchanged). Conv2d has issue when act_block_h_override is used. Assert that act_block_h_override in tiles needs be a divisor of per_core_out_matrix_height_ntiles was wrongly disabled. In turn each core would produce more output than than it was allocated and override random data. which is not visible in op unit tests. This change regress unet shallow device perf. 1122 -> 1053. Most of the regressino comes from a workaround to make model work determinically which converts one sharded tensor to l1 interleaved tensor. Root cause is unknow but we did hit this issue before and it went away, not it's back.
pavlejosipovic
force-pushed
the
pjosipovic/fix_cb_asserts
branch
from
a9f2805 to
aa0d30e
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.
Halo op has a problem where manipulating
tt::tt_metal::LegacyShape in compute_output_specs
is not doing the expected thing (value is unchanged).
Conv2d has issue when act_block_h_override is used. Assert that act_block_h_override in tiles needs
be a divisor of per_core_out_matrix_height_ntiles
was wrongly disabled.
In turn each core would produce more output than
than it was allocated and override random data.
which is not visible in op unit tests.
This change regress unet shallow device perf.
1122 -> 1053.
Most of the regressino comes from a workaround
to make model work determinically which converts
one sharded tensor to l1 interleaved tensor.
Root cause is unknow but we did hit this issue
before and it went away, not it's back.
Checklist