#0: Update get_semaphore to return 16B aligned semaphore addresses - …

…fix bug where host and dev semaphore addresses did not match and simplify semaphore constants
tenstorrent · Jan 19, 2024 · e08d60e · e08d60e
1 parent fa0dd3c
commit e08d60e
Show file tree

Hide file tree

Showing 5 changed files with 8 additions and 10 deletions.
diff --git a/tests/tt_metal/tt_metal/test_core_range_set.cpp b/tests/tt_metal/tt_metal/test_core_range_set.cpp
@@ -53,7 +53,7 @@ void check_semaphores_are_initialized(tt_metal::Device *device, const CoreRangeS
                 std::vector<uint32_t> res;
                 tt_metal::detail::ReadFromDeviceL1(device, logical_core, SEMAPHORE_BASE, SEMAPHORE_SIZE, res);
                 std::vector<uint32_t> filtered_res;
-                constexpr static uint32_t num_u32_to_skip = UINT32_WORDS_PER_SEMAPHORE * sizeof(uint32_t);
+                constexpr static uint32_t num_u32_to_skip = sizeof(uint32_t);
                 for (int i = 0; i < res.size(); i+=num_u32_to_skip) {
                     filtered_res.push_back(res.at(i));
                 }

diff --git a/tests/tt_metal/tt_metal/unit_tests/basic/initialize_semaphores.cpp b/tests/tt_metal/tt_metal/unit_tests/basic/initialize_semaphores.cpp
@@ -68,7 +68,7 @@ void create_and_read_max_num_semaphores(
         uint32_t initial_value = i;
         auto semaphore_addr = tt_metal::CreateSemaphore(program, core_range, initial_value);
         golden.push_back(initial_value);
-        ASSERT_TRUE(semaphore_addr == SEMAPHORE_BASE + (ALIGNED_SIZE_PER_SEMAPHORE * i));
+        ASSERT_TRUE(semaphore_addr == SEMAPHORE_BASE + (L1_ALIGNMENT * i));
     }
 
     ASSERT_TRUE(tt_metal::detail::ConfigureDeviceWithProgram(device, program));
@@ -79,8 +79,8 @@ void create_and_read_max_num_semaphores(
             std::vector<uint32_t> res;
             for (uint32_t i = 0; i < NUM_SEMAPHORES; i++) {
                 std::vector<uint32_t> single_val;
-                uint32_t semaphore_addr = SEMAPHORE_BASE + (ALIGNED_SIZE_PER_SEMAPHORE * i);
-                uint32_t semaphore_size = UINT32_WORDS_PER_SEMAPHORE * sizeof(uint32_t);
+                uint32_t semaphore_addr = SEMAPHORE_BASE + (L1_ALIGNMENT * i);
+                uint32_t semaphore_size = sizeof(uint32_t);
                 tt_metal::detail::ReadFromDeviceL1(device, logical_core, semaphore_addr, semaphore_size, single_val);
                 ASSERT_TRUE(single_val.size() == 1);
                 res.push_back(single_val.at(0));

diff --git a/tt_metal/hostdevcommon/common_runtime_address_map.h b/tt_metal/hostdevcommon/common_runtime_address_map.h
@@ -29,12 +29,10 @@ constexpr static std::uint32_t NUM_CIRCULAR_BUFFERS = 32;
 constexpr static std::uint32_t UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG = 4;
 constexpr static std::uint32_t CIRCULAR_BUFFER_CONFIG_SIZE = NUM_CIRCULAR_BUFFERS * UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG * sizeof(uint32_t);
 
-// 4 semaphores per core aligned to 16B
+// 4 uint32_t semaphores per core aligned to 16B
 constexpr static std::uint32_t SEMAPHORE_BASE = CIRCULAR_BUFFER_CONFIG_BASE + CIRCULAR_BUFFER_CONFIG_SIZE;
 constexpr static std::uint32_t NUM_SEMAPHORES = 4;
-constexpr static std::uint32_t UINT32_WORDS_PER_SEMAPHORE = 1;
-constexpr static std::uint32_t ALIGNED_SIZE_PER_SEMAPHORE = (((UINT32_WORDS_PER_SEMAPHORE * sizeof(uint32_t)) + L1_ALIGNMENT - 1) / L1_ALIGNMENT) * L1_ALIGNMENT;
-constexpr static std::uint32_t SEMAPHORE_SIZE = NUM_SEMAPHORES * ALIGNED_SIZE_PER_SEMAPHORE;
+constexpr static std::uint32_t SEMAPHORE_SIZE = NUM_SEMAPHORES * L1_ALIGNMENT;
 
 // Debug printer buffers - A total of 5*PRINT_BUFFER_SIZE starting at PRINT_BUFFER_NC address
 constexpr static std::uint32_t PRINT_BUFFER_SIZE = 204; // per thread

diff --git a/tt_metal/hw/inc/dataflow_api.h b/tt_metal/hw/inc/dataflow_api.h
@@ -1148,7 +1148,7 @@ FORCE_INLINE void noc_async_write_tile(
 
 FORCE_INLINE
 uint32_t get_semaphore(uint32_t semaphore_id) {
-    return SEMAPHORE_BASE + semaphore_id * sizeof(uint32_t);
+    return SEMAPHORE_BASE + semaphore_id * L1_ALIGNMENT;
 }
 
 FORCE_INLINE

diff --git a/tt_metal/tt_metal.cpp b/tt_metal/tt_metal.cpp
@@ -55,7 +55,7 @@ std::optional<uint32_t> get_semaphore_address(const Program &program, const Core
             }
             uint32_t addr = num_semaphores == 0
                                 ? SEMAPHORE_BASE
-                                : program.semaphore_address(num_semaphores - 1) + ALIGNED_SIZE_PER_SEMAPHORE;
+                                : program.semaphore_address(num_semaphores - 1) + L1_ALIGNMENT;
             if (!address.has_value()) {
                 address = addr;
             } else if (addr != address) {