#13609: Uplift dram and l1 allocators to use dram/l1 specific alignment #13762

abhullar-tt · 2024-10-11T21:17:02Z

Ticket

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

…ixed

…rm_with_block_sharded_v2_8x8_grid

tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h

ttnn/cpp/ttnn/operations/normalization/groupnorm/device/multi_core/groupnorm_op_multi_core.cpp

abhullar-tt · 2025-01-06T17:25:09Z

@llongTT could you re-run post-commit, BH post commit, Model tests, and ttnn sweeps and update the links in the PR description

yugaoTT · 2025-01-06T17:25:21Z

Please also run the device perf and model pipelines, there are CNN models using group norm so need to make sure those are passing

abhullar-tt · 2025-01-06T17:26:46Z

@cmaryanTT we may need to announce when this change goes in because using l1/dram specific alignment for allocating buffers could expose noc violations when doing misaligned dram -> l1 reads. We could explain how ops are handling this

tt_metal/impl/dispatch/memcpy.hpp

…name for l1 alignment

…lignment(dst_buffer->buffer_type()) to dst_buffer->alignment()

llongTT · 2025-01-06T18:30:54Z

@llongTT could you re-run post-commit, BH post commit, Model tests, and ttnn sweeps and update the links in the PR description

Have them running now.

llongTT · 2025-01-06T18:31:20Z

Please also run the device perf and model pipelines, there are CNN models using group norm so need to make sure those are passing

Have them running now.

cmaryanTT · 2025-01-06T18:35:39Z

@cmaryanTT we may need to announce when this change goes in because using l1/dram specific alignment for allocating buffers could expose noc violations when doing misaligned dram -> l1 reads. We could explain how ops are handling this

Noted. @llongTT - please announce on #tt-metal-developers with a brief description of the change and the risks.

Also, @llongTT will be presenting this in the learning forum in a couple of weeks.

…movement from DRAM to L1

github-actions

⚠️ Clang-Tidy found issue(s) with the introduced code (1/1)

github-actions · 2025-01-09T23:13:26Z

tt_metal/impl/buffers/buffer.cpp

@@ -449,7 +449,7 @@ DeviceAddr Buffer::bank_local_page_address(uint32_t bank_id, uint32_t page_index
 }

 uint32_t Buffer::alignment() const {
-    return this->allocator_->config.alignment;
+    return this->device_->get_allocator_alignment(this->buffer_type_);


⚠️ clang-diagnostic-error ⚠️
no matching member function for call to get_allocator_alignment

github-actions · 2025-01-09T23:13:27Z

tt_metal/impl/device/device.cpp

@@ -1406,14 +1407,14 @@ allocator::Statistics Device::get_memory_allocation_statistics(const BufferType
    return allocator::get_statistics(*allocator, buffer_type);
 }

-uint32_t Device::get_allocator_alignment() const {
+uint32_t Device::get_allocator_alignment(const BufferType &buffer_type) const {


⚠️ clang-diagnostic-error ⚠️
out-of-line definition of get_allocator_alignment does not match any declaration in tt::tt_metal::Device

github-actions · 2025-01-09T23:13:27Z

tt_metal/impl/device/device.hpp

@@ -129,8 +129,8 @@ class Device : public IDevice {
    allocator::Statistics get_memory_allocation_statistics(const BufferType &buffer_type) const override;
    allocator::Statistics get_memory_allocation_statistics(const BufferType &buffer_type, SubDeviceId sub_device_id) const override;

-    uint32_t get_allocator_alignment() const override;
-    uint32_t get_allocator_alignment(SubDeviceId sub_device_id) const override;
+    uint32_t get_allocator_alignment(const BufferType &buffer_type) const override;


⚠️ clang-diagnostic-error ⚠️
non-virtual member function marked override hides virtual member functions

github-actions · 2025-01-09T23:13:27Z

tt_metal/impl/device/device.hpp

-    uint32_t get_allocator_alignment() const override;
-    uint32_t get_allocator_alignment(SubDeviceId sub_device_id) const override;
+    uint32_t get_allocator_alignment(const BufferType &buffer_type) const override;
+    uint32_t get_allocator_alignment(const BufferType &buffer_type, SubDeviceId sub_device_id) const override;


⚠️ clang-diagnostic-error ⚠️
non-virtual member function marked override hides virtual member functions

github-actions · 2025-01-09T23:13:27Z

tests/tt_metal/tt_metal/api/allocator/test_l1_banking_allocator.cpp

@@ -19,7 +19,7 @@ uint64_t get_alloc_limit(const tt::tt_metal::IDevice* device) {
    auto dispatch_core_config = dispatch_core_manager::instance().get_dispatch_core_config(device->id());
    auto storage_core_bank_size =
        tt::get_storage_core_bank_size(device->id(), device->num_hw_cqs(), dispatch_core_config);
-    const uint32_t allocator_alignment = device->get_allocator_alignment();
+    const uint32_t allocator_alignment = device->get_allocator_alignment(BufferType::L1);


⚠️ clang-diagnostic-error ⚠️
no matching member function for call to get_allocator_alignment

github-actions · 2025-01-09T23:13:27Z

tt_metal/impl/program/program.cpp

@@ -788,7 +788,7 @@ void detail::Program_::allocate_circular_buffers(const IDevice* device) {
                }
            }
        }
-        computed_addr = align(computed_addr, device->get_allocator_alignment());
+        computed_addr = align(computed_addr, device->get_allocator_alignment(BufferType::DRAM));


⚠️ clang-diagnostic-error ⚠️
no matching member function for call to get_allocator_alignment

github-actions

⚠️ Clang-Tidy found issue(s) with the introduced code (1/1)

ttnn/cpp/ttnn/operations/data_movement/transpose/device/transpose_program_factory.cpp

github-actions · 2025-01-10T21:23:21Z

ttnn/cpp/ttnn/operations/data_movement/transpose/device/transpose_program_factory.cpp

@@ -676,7 +676,6 @@
    // 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 ? hal::get_dram_alignment()
-                                                                                     : hal::get_l1_alignment();
    bool misaligned = alignment > sub_tile_line_bytes;


⚠️ clang-diagnostic-error ⚠️
expected ( for function-style cast or type construction

…ose_program_factory.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>

github-actions

⚠️ Clang-Tidy found issue(s) with the introduced code (1/1)

github-actions · 2025-01-10T21:48:39Z

ttnn/cpp/ttnn/operations/data_movement/transpose/device/transpose_program_factory.cpp

@@ -675,8 +675,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 ? hal::get_dram_alignment()
-                                                                                     : hal::get_l1_alignment();
+    uint32_t alignment = dst_buffer->buffer_type() == tt::tt_metal::BufferType::DRAM ? hal::get_dram_alignment();


⚠️ clang-diagnostic-error ⚠️
expected :

Suggested change

uint32_t alignment = dst_buffer->buffer_type() == tt::tt_metal::BufferType::DRAM ? hal::get_dram_alignment();

uint32_t alignment = dst_buffer->buffer_type() == tt::tt_metal::BufferType::DRAM ? hal::get_dram_alignment(): ;

github-actions · 2025-01-10T21:48:39Z

ttnn/cpp/ttnn/operations/data_movement/transpose/device/transpose_program_factory.cpp

@@ -675,8 +675,7 @@
    // 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 ? hal::get_dram_alignment()
-                                                                                     : hal::get_l1_alignment();
+    uint32_t alignment = dst_buffer->buffer_type() == tt::tt_metal::BufferType::DRAM ? hal::get_dram_alignment();


⚠️ clang-diagnostic-error ⚠️
expected expression

abhullar-tt mentioned this pull request Oct 11, 2024

allocator uses 32B alignment for both DRAM and L1 #13609

Open

abhullar-tt temporarily deployed to dev October 11, 2024 21:18 — with GitHub Actions Inactive

abhullar-tt temporarily deployed to dev October 11, 2024 21:19 — with GitHub Actions Inactive

abhullar-tt temporarily deployed to dev October 11, 2024 21:20 — with GitHub Actions Inactive

abhullar-tt temporarily deployed to dev October 11, 2024 21:22 — with GitHub Actions Inactive

abhullar-tt linked an issue Oct 11, 2024 that may be closed by this pull request

allocator uses 32B alignment for both DRAM and L1 #13609

Open

abhullar-tt had a problem deploying to dev October 11, 2024 21:29 — with GitHub Actions Failure

abhullar-tt temporarily deployed to dev October 11, 2024 21:29 — with GitHub Actions Inactive

abhullar-tt temporarily deployed to dev October 11, 2024 21:32 — with GitHub Actions Inactive

abhullar-tt temporarily deployed to dev October 11, 2024 21:35 — with GitHub Actions Inactive

llongTT added 2 commits January 3, 2025 19:00

#13609: enable test_permute_5d_blocked as the memory issue has been f…

fead9bc

…ixed

#13609: explicitely pack l1 for group norm unit test of test_group_no…

ffbf4a4

…rm_with_block_sharded_v2_8x8_grid

sjameelTT reviewed Jan 6, 2025

View reviewed changes

tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h Show resolved Hide resolved

ttnn/cpp/ttnn/operations/normalization/groupnorm/device/multi_core/groupnorm_op_multi_core.cpp Outdated Show resolved Hide resolved

jvegaTT reviewed Jan 6, 2025

View reviewed changes

tt_metal/impl/dispatch/memcpy.hpp Show resolved Hide resolved

jvegaTT approved these changes Jan 6, 2025

View reviewed changes

llongTT added 4 commits January 6, 2025 17:49

#13609: address some review comment, regarding comments and variable …

45ed294

…name for l1 alignment

#13609: simplify alignment function call from device->get_allocator_a…

436d7e8

…lignment(dst_buffer->buffer_type()) to dst_buffer->alignment()

#13609: switch from round_up_to_mul16 to tt::round_up

6cca1b9

Merge branch 'main' into abhullar/diff-aligns

9964d75

sjameelTT approved these changes Jan 6, 2025

View reviewed changes

llongTT self-assigned this Jan 6, 2025

llongTT added 2 commits January 9, 2025 22:56

#13609: set L1 bank address to be DRAM aligned, to facilitate memory …

a977eb4

…movement from DRAM to L1

Merge branch 'main' into abhullar/diff-aligns

69cd54d

github-actions bot reviewed Jan 9, 2025

View reviewed changes

llongTT added 4 commits January 10, 2025 00:42

#13609: fix the compile error for non-virtual function

7892dbf

#13609: set the device->get_allocator_alignment to purely virtual

3121bcd

Merge branch 'main' into abhullar/diff-aligns

ff6b6e0

Merge branch 'main' into abhullar/diff-aligns

1170b01

github-actions bot reviewed Jan 10, 2025

View reviewed changes

llongTT and others added 2 commits January 10, 2025 13:33

Update ttnn/cpp/ttnn/operations/data_movement/transpose/device/transp…

3544064

…ose_program_factory.cpp Co-authored-by: github-actions[bot] <41898282+github-actions[bot]@users.noreply.github.com>

#13609: Fix a missing semicolon to pass the build

0a4aae2

github-actions bot reviewed Jan 10, 2025

View reviewed changes

llongTT added 2 commits January 10, 2025 22:11

#13609: mis-copied line from github merge

8327eca

Merge branch 'main' into abhullar/diff-aligns

ec1c03e

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

#13609: Uplift dram and l1 allocators to use dram/l1 specific alignment #13762

#13609: Uplift dram and l1 allocators to use dram/l1 specific alignment #13762

abhullar-tt commented Oct 11, 2024 •

edited by llongTT

Loading

abhullar-tt commented Jan 6, 2025

yugaoTT commented Jan 6, 2025

abhullar-tt commented Jan 6, 2025

llongTT commented Jan 6, 2025

llongTT commented Jan 6, 2025

cmaryanTT commented Jan 6, 2025

github-actions bot left a comment

github-actions bot Jan 9, 2025

github-actions bot Jan 9, 2025

github-actions bot Jan 9, 2025

github-actions bot Jan 9, 2025

github-actions bot Jan 9, 2025

github-actions bot Jan 9, 2025

github-actions bot left a comment

github-actions bot Jan 10, 2025

github-actions bot left a comment

github-actions bot Jan 10, 2025

github-actions bot Jan 10, 2025

	uint32_t alignment = dst_buffer->buffer_type() == tt::tt_metal::BufferType::DRAM ? hal::get_dram_alignment();
	uint32_t alignment = dst_buffer->buffer_type() == tt::tt_metal::BufferType::DRAM ? hal::get_dram_alignment(): ;

#13609: Uplift dram and l1 allocators to use dram/l1 specific alignment #13762

Are you sure you want to change the base?

#13609: Uplift dram and l1 allocators to use dram/l1 specific alignment #13762

Conversation

abhullar-tt commented Oct 11, 2024 • edited by llongTT Loading

Ticket

Problem description

What's changed

Checklist

abhullar-tt commented Jan 6, 2025

yugaoTT commented Jan 6, 2025

abhullar-tt commented Jan 6, 2025

llongTT commented Jan 6, 2025

llongTT commented Jan 6, 2025

cmaryanTT commented Jan 6, 2025

github-actions bot left a comment

Choose a reason for hiding this comment

github-actions bot Jan 9, 2025

Choose a reason for hiding this comment

github-actions bot Jan 9, 2025

Choose a reason for hiding this comment

github-actions bot Jan 9, 2025

Choose a reason for hiding this comment

github-actions bot Jan 9, 2025

Choose a reason for hiding this comment

github-actions bot Jan 9, 2025

Choose a reason for hiding this comment

github-actions bot Jan 9, 2025

Choose a reason for hiding this comment

github-actions bot left a comment

Choose a reason for hiding this comment

github-actions bot Jan 10, 2025

Choose a reason for hiding this comment

github-actions bot left a comment

Choose a reason for hiding this comment

github-actions bot Jan 10, 2025

Choose a reason for hiding this comment

github-actions bot Jan 10, 2025

Choose a reason for hiding this comment

abhullar-tt commented Oct 11, 2024 •

edited by llongTT

Loading