Migrate CBs and Sems to kernel config ring buffer

tests/tt_metal/tt_metal/test_core_range_set.cpp

@@ -46,18 +46,22 @@ void check_program_is_mapped_to_correct_cores(const tt_metal::Program &program,
     }
 }
 
-void check_semaphores_are_initialized(tt_metal::Device *device, const CoreRangeSet &core_range_set, const std::vector<uint32_t> &golden_sem_values) {
+void check_semaphores_are_initialized(tt_metal::Device *device, tt_metal::Program& program, const CoreRangeSet &core_range_set, const std::vector<uint32_t> &golden_sem_values) {
     for (auto core_range : core_range_set.ranges()) {
         for (auto x = core_range.start_coord.x; x <= core_range.end_coord.x; x++) {
             for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                 auto logical_core = CoreCoord{x, y};
                 std::vector<uint32_t> res;
-                tt_metal::detail::ReadFromDeviceL1(device, logical_core, SEMAPHORE_BASE, SEMAPHORE_SIZE, res);
+                tt_metal::detail::ReadFromDeviceL1(device, logical_core,
+                    program.get_sem_base_addr(device, logical_core, CoreType::WORKER),
+                    program.get_sem_size(device, logical_core, CoreType::WORKER),
+                    res);
                 std::vector<uint32_t> filtered_res;
-                constexpr static uint32_t num_u32_to_skip = sizeof(uint32_t);
+                constexpr static uint32_t num_u32_to_skip = L1_ALIGNMENT / sizeof(uint32_t);
                 for (int i = 0; i < res.size(); i+=num_u32_to_skip) {
                     filtered_res.push_back(res.at(i));
                 }
+
                 TT_FATAL(filtered_res == golden_sem_values);
             }
         }
@@ -142,9 +146,8 @@ bool test_program_specified_with_core_range_set(tt_metal::Device *device, tt_met
         tt_metal::ComputeConfig{.compile_args = compute_kernel_args}
     );
 
-    auto size_per_semaphore = SEMAPHORE_SIZE / NUM_SEMAPHORES;
     std::vector<uint32_t> golden_sem_values;
-    for (uint32_t i = 0; i < NUM_SEMAPHORES; i++) {
+    for (uint32_t i = 0; i < tt_metal::NUM_SEMAPHORES; i++) {
         uint32_t initial_value = i;
         tt_metal::CreateSemaphore(program, core_range_set, initial_value);
         golden_sem_values.push_back(initial_value);
@@ -154,10 +157,6 @@ bool test_program_specified_with_core_range_set(tt_metal::Device *device, tt_met
 
     tt_metal::detail::CompileProgram(device, program);
 
-    pass &= tt_metal::detail::ConfigureDeviceWithProgram(device, program);
-
-    check_semaphores_are_initialized(device, core_range_set, golden_sem_values);
-
     std::vector<uint32_t> src_vec = create_random_vector_of_bfloat16(
             buffer_size, 100, std::chrono::system_clock::now().time_since_epoch().count());
     tt_metal::detail::WriteToBuffer(src_dram_buffer, src_vec);
@@ -191,6 +190,8 @@ bool test_program_specified_with_core_range_set(tt_metal::Device *device, tt_met
 
     tt_metal::detail::LaunchProgram(device, program);
 
+    check_semaphores_are_initialized(device, program, core_range_set, golden_sem_values);
+
     for (const auto &[core, dst_l1_buffer] : core_to_l1_buffer) {
         std::vector<uint32_t> result_vec;
         tt_metal::detail::ReadFromBuffer(dst_l1_buffer, result_vec);

tests/tt_metal/tt_metal/test_kernels/compute/3T/matmul_large_block_zm/zm_3m_pack.cpp

@@ -86,7 +86,6 @@ void pack_main()
 {
 uint32_t in0_block_w = get_compile_time_arg_val(0);
 llk_pack_init();
-llk_setup_outputs();
 llk_pack_dest_init<DstTileFaceLayout::RowMajor, false>();
 llk_init_packer_dest_offset_registers<DstTileFaceLayout::RowMajor,false>();
 llk_pack_hw_configure_disaggregated<false>(16);

tests/tt_metal/tt_metal/test_kernels/compute/3T/matmul_large_block_zm/zm_3m_unpack.cpp

@@ -122,7 +122,6 @@ inline void unpack_for_matmul_output_row(
 void unpack_main()
 {
 uint32_t in0_block_w = get_compile_time_arg_val(0);
-llk_setup_operands();
 llk_unpack_AB_matmul_init(0);
 // inner block size in tiles
 uint32_t in0_num_subblocks = get_compile_time_arg_val(1);

tests/tt_metal/tt_metal/test_kernels/compute/3T/untilize_A_and_eltwise_binary/chlkc_pack.cpp

@@ -15,7 +15,6 @@ uint32_t per_core_block_r_tiles = get_compile_time_arg_val(1);
 uint32_t per_core_block_c_tiles = get_compile_time_arg_val(2);
 llk_pack_init();
 llk_pack_hw_configure_disaggregated<false>(16);
-llk_setup_outputs();
 llk_pack_dest_init<DstTileFaceLayout::RowMajor, false>();
 
 for (uint32_t block = 0; block < per_core_num_blocks; block++) {

tests/tt_metal/tt_metal/test_kernels/compute/3T/untilize_A_and_eltwise_binary/chlkc_unpack.cpp

@@ -16,7 +16,6 @@ uint32_t per_core_num_blocks = get_compile_time_arg_val(0);
 uint32_t per_core_block_r_tiles = get_compile_time_arg_val(1);
 uint32_t per_core_block_c_tiles = get_compile_time_arg_val(2);
 
-llk_setup_operands();
 llk_unpack_AB_hw_configure_disaggregated<BroadcastType::NONE>(0,1);
 // llk_unpack_untilize_hw_configure_disaggregated(0);
 

tests/tt_metal/tt_metal/test_kernels/compute/eltwise_copy_3m.cpp

@@ -39,7 +39,6 @@ void pack_main()
     int __outer_loop_iter;
     llk_pack_init();
     llk_pack_hw_configure_disaggregated<false>(16);
-    llk_setup_outputs();
     llk_pack_dest_init<DstTileFaceLayout::RowMajor, false>();
     constexpr uint32_t per_core_tile_cnt = get_compile_time_arg_val(0);
     for (uint32_t b = 0; b < per_core_tile_cnt; ++b) {
@@ -56,7 +55,6 @@ void pack_main()
 void unpack_main()
 {
     int __outer_loop_iter;
-    llk_setup_operands();
     UNPACK(( llk_unpack_A_init<BroadcastType::NONE, false, EltwiseBinaryReuseDestType::NONE>()  ));
     UNPACK(( llk_unpack_A_hw_configure_disaggregated<BroadcastType::NONE>(0) ));
     constexpr uint32_t per_core_tile_cnt = get_compile_time_arg_val(0);

tests/tt_metal/tt_metal/test_kernels/compute/untilA_elwbin_3m.cpp

@@ -59,7 +59,6 @@ uint32_t per_core_num_blocks = get_compile_time_arg_val(0);
 uint32_t per_core_block_r_tiles = get_compile_time_arg_val(1);
 uint32_t per_core_block_c_tiles = get_compile_time_arg_val(2);
 
-llk_setup_operands();
 llk_unpack_AB_hw_configure_disaggregated<BroadcastType::NONE>(0,1);
 // llk_unpack_untilize_hw_configure_disaggregated(0);
 
@@ -99,7 +98,6 @@ void pack_main()
     uint32_t per_core_block_c_tiles = get_compile_time_arg_val(2);
     llk_pack_init();
     llk_pack_hw_configure_disaggregated<false>(16);
-    llk_setup_outputs();
     llk_pack_dest_init<DstTileFaceLayout::RowMajor, false>();
 
     for (uint32_t block = 0; block < per_core_num_blocks; block++) {

tests/tt_metal/tt_metal/test_kernels/dataflow/unit_tests/command_queue/random_program.cpp

@@ -39,22 +39,28 @@ void MAIN  {
     constexpr volatile uint32_t page_size = get_compile_time_arg_val(7);
 
     for (uint32_t i = 0; i < num_cbs; i++) {
-        uint32_t cb_val = reinterpret_cast<volatile tt_l1_ptr uint32_t*>(CIRCULAR_BUFFER_CONFIG_BASE + i * 16)[3];
+        tt_l1_ptr mailboxes_t* const mailboxes = (tt_l1_ptr mailboxes_t*)(MEM_MAILBOX_BASE);
+        uint32_t kernel_config_base = mailboxes->launch.kernel_config.kernel_config_base;
+        uint32_t tt_l1_ptr *cb_l1_base = (uint32_t tt_l1_ptr *)(kernel_config_base +
+            mailboxes->launch.kernel_config.cb_offset);
+        uint32_t cb_val = reinterpret_cast<volatile tt_l1_ptr uint32_t*>(cb_l1_base + i * 4)[3];
         uint32_t expected = ((i + 1) * page_size) >> 4;
         if (cb_val != expected) {
             DPRINT << "Problem with CB idx: " << i << " Expected: " << expected << " Got: " << cb_val << ENDL();
             while(true); // Purposefully hang the kernel if CBs did not arrive correctly
         }
     }
 
+#ifdef DATA_MOVEMENT
     for (uint32_t i = 0; i < num_sems; i++) {
-        uint32_t sem_val = reinterpret_cast<volatile tt_l1_ptr uint32_t*>(SEMAPHORE_BASE + i * 16)[0];
+        uint32_t sem_val = reinterpret_cast<volatile tt_l1_ptr uint32_t*>(get_semaphore(i))[0];
         uint32_t expected = i + 1;
         if (sem_val != expected) {
             DPRINT << "Problem with Sem idx: " << i << " Expected: " << expected << " Got: " << sem_val << ENDL();
             while(true); // Purposefully hang the kernel if semaphores did not arrive correctly
         }
     }
+#endif
 
     for (uint32_t i = 0; i < num_unique_rt_args; i++) {
         uint32_t rt_arg = get_arg_val<uint32_t>(i);

tests/tt_metal/tt_metal/unit_tests/basic/initialize_semaphores.cpp

@@ -79,7 +79,7 @@ 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 + (L1_ALIGNMENT * i);
+                uint32_t semaphore_addr = program.get_sem_base_addr(device, logical_core, CoreType::WORKER) + (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);

tests/tt_metal/tt_metal/unit_tests/circular_buffer/test_CircularBuffer_allocation.cpp

@@ -25,7 +25,7 @@ void validate_cb_address(Program &program, Device *device, const CoreRangeSet &c
             for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                 CoreCoord core_coord(x, y);
                 tt::tt_metal::detail::ReadFromDeviceL1(
-                    device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                    device, core_coord, program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                 std::map<uint8_t, uint32_t> address_per_buffer_index = core_to_address_per_buffer_index.at(core_coord);
 
@@ -284,6 +284,7 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferAddress) {
 
 TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
   for (unsigned int id = 0; id < num_devices_; id++) {
+    Device *device = this->devices_.at(id);
     Program program;
     CBConfig cb_config;
     CoreCoord core0(0, 0);
@@ -306,7 +307,7 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
         expected_cb_addr += cb_config.page_size;
     }
 
-    detail::LaunchProgram(this->devices_.at(id), program);
+    detail::LaunchProgram(device, program);
 
     vector<uint32_t> cb_config_vector;
     uint32_t cb_config_buffer_size = NUM_CIRCULAR_BUFFERS * UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG * sizeof(uint32_t);
@@ -316,7 +317,7 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
             for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                 CoreCoord core_coord(x, y);
                 tt::tt_metal::detail::ReadFromDeviceL1(
-                    this->devices_.at(id), core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                    device, core_coord, program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                 std::map<uint8_t, uint32_t> address_per_buffer_index = golden_addresses_per_core.at(core_coord);
                 std::map<uint8_t, uint32_t> num_pages_per_buffer_index = golden_num_pages_per_core.at(core_coord);
@@ -333,15 +334,15 @@ TEST_F(DeviceFixture, TestUpdateCircularBufferPageSize) {
     UpdateCircularBufferPageSize(program, cb_ids[1], 1, cb_config.page_size / 2);
     golden_num_pages_per_core[core0][1] = 2;
 
-    detail::LaunchProgram(this->devices_.at(id), program);
+    detail::LaunchProgram(device, program);
 
     // addresses should not be changed
     for (const CoreRange &core_range : cr_set.ranges()) {
         for (auto x = core_range.start_coord.x; x <= core_range.end_coord.x; x++) {
             for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                 CoreCoord core_coord(x, y);
                 tt::tt_metal::detail::ReadFromDeviceL1(
-                    this->devices_.at(id), core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                    device, core_coord, program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                 std::map<uint8_t, uint32_t> address_per_buffer_index = golden_addresses_per_core.at(core_coord);
                 std::map<uint8_t, uint32_t> num_pages_per_buffer_index = golden_num_pages_per_core.at(core_coord);

tests/tt_metal/tt_metal/unit_tests/circular_buffer/test_CircularBuffer_creation.cpp

@@ -28,7 +28,7 @@ bool test_cb_config_written_to_core(Program &program, Device *device, const Core
                 for (auto y = core_range.start_coord.y; y <= core_range.end_coord.y; y++) {
                     CoreCoord core_coord(x, y);
                     tt::tt_metal::detail::ReadFromDeviceL1(
-                        device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+                        device, core_coord, program.get_sem_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
 
                     for (const auto &[buffer_index, golden_cb_config] : cb_config_per_buffer_index) {
                         auto base_index = UINT32_WORDS_PER_CIRCULAR_BUFFER_CONFIG * buffer_index;

tests/tt_metal/tt_metal/unit_tests_fast_dispatch/command_queue/test_EnqueueProgram.cpp

@@ -76,7 +76,7 @@ std::vector<CBHandle> initialize_dummy_circular_buffers(Program& program, const
     return cb_handles;
 }
 
-bool cb_config_successful(Device* device, const DummyProgramMultiCBConfig & program_config){
+bool cb_config_successful(Device* device, Program &program, const DummyProgramMultiCBConfig & program_config){
     bool pass = true;
 
     // Need to use old APIs to read since we cannot allocate a buffer in the reserved space we're trying
@@ -86,7 +86,9 @@ bool cb_config_successful(Device* device, const DummyProgramMultiCBConfig & prog
 
     for (const CoreRange& core_range : program_config.cr_set.ranges()) {
         for (const CoreCoord& core_coord : core_range) {
-            tt::tt_metal::detail::ReadFromDeviceL1(device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+            tt::tt_metal::detail::ReadFromDeviceL1(device, core_coord,
+                program.get_sem_base_addr(device, core_coord, CoreType::WORKER),
+                cb_config_buffer_size, cb_config_vector);
 
             uint32_t cb_addr = L1_UNRESERVED_BASE;
             for (uint32_t i = 0; i < program_config.cb_config_vector.size(); i++) {
@@ -115,7 +117,7 @@ bool test_dummy_EnqueueProgram_with_cbs(Device* device, CommandQueue& cq, DummyP
     EnqueueProgram(cq, program, is_blocking_op);
     Finish(cq);
 
-    return cb_config_successful(device, program_config);
+    return cb_config_successful(device, program, program_config);
 }
 
 bool test_dummy_EnqueueProgram_with_cbs_update_size(Device* device, CommandQueue& cq, const DummyProgramMultiCBConfig& program_config) {
@@ -126,7 +128,7 @@ bool test_dummy_EnqueueProgram_with_cbs_update_size(Device* device, CommandQueue
     EnqueueProgram(cq, program, false);
     Finish(cq);
 
-    const bool is_cb_config_before_update_successful = cb_config_successful(device, program_config);
+    const bool is_cb_config_before_update_successful = cb_config_successful(device, program, program_config);
 
     DummyProgramMultiCBConfig program_config_2 = program_config;
     for (uint32_t cb_id = 0; cb_id < program_config.cb_config_vector.size(); cb_id++) {
@@ -139,7 +141,7 @@ bool test_dummy_EnqueueProgram_with_cbs_update_size(Device* device, CommandQueue
     EnqueueProgram(cq, program, false);
     Finish(cq);
 
-    const bool is_cb_config_after_update_successful = cb_config_successful(device, program_config_2);
+    const bool is_cb_config_after_update_successful = cb_config_successful(device, program, program_config_2);
     return is_cb_config_before_update_successful && is_cb_config_after_update_successful;
 }
 
@@ -163,7 +165,8 @@ bool test_dummy_EnqueueProgram_with_sems(Device* device, CommandQueue& cq, Progr
             vector<uint32_t> semaphore_vals;
             uint32_t expected_semaphore_vals_for_core_idx = 0;
             const uint32_t semaphore_buffer_size = program_config.num_sems * L1_ALIGNMENT;
-            tt::tt_metal::detail::ReadFromDeviceL1(device, core_coord, SEMAPHORE_BASE, semaphore_buffer_size, semaphore_vals);
+            uint32_t semaphore_base = program.get_sem_base_addr(device, core_coord, CoreType::WORKER);
+            tt::tt_metal::detail::ReadFromDeviceL1(device, core_coord, semaphore_base, semaphore_buffer_size, semaphore_vals);
             for (uint32_t i = 0; i < semaphore_vals.size(); i += (L1_ALIGNMENT / sizeof(uint32_t)))
             {
                 const bool is_semaphore_value_correct = semaphore_vals[i] == expected_semaphore_vals_for_core[expected_semaphore_vals_for_core_idx];
@@ -775,7 +778,8 @@ TEST_F(CommandQueueSingleCardFixture, TestMultiCBSharedAddressSpaceSentSingleCor
         vector<uint32_t> cb_config_vector;
 
         tt::tt_metal::detail::ReadFromDeviceL1(
-            device, core_coord, CIRCULAR_BUFFER_CONFIG_BASE, cb_config_buffer_size, cb_config_vector);
+            device, core_coord,
+            program.get_cb_base_addr(device, core_coord, CoreType::WORKER), cb_config_buffer_size, cb_config_vector);
         uint32_t cb_addr = L1_UNRESERVED_BASE;
         uint32_t intermediate_index = intermediate_cb * sizeof(uint32_t);
 

tt_metal/hostdevcommon/common_runtime_address_map.h

@@ -51,33 +51,22 @@ constexpr static std::uint32_t PROFILER_RISC_COUNT = 5;
 static_assert (PROFILER_FULL_HOST_BUFFER_SIZE_PER_RISC > PROFILER_L1_BUFFER_SIZE);
 
 // Kernel config buffer is WIP
-// Will eventually move CBs/Sems and likely kernel bins into this buffer
-// Size is presently based on the old size of the RTAs (large enough to hold 1 set)
+// Size is presently based on the old sizes of the RTAs + CB config + Sems
 // plus some extra space freed up in the mem map
 constexpr static std::uint32_t L1_KERNEL_CONFIG_BASE = PROFILER_L1_END_ADDRESS;
-constexpr static std::uint32_t L1_KERNEL_CONFIG_SIZE = 4 * 1024 + 256;
+constexpr static std::uint32_t L1_KERNEL_CONFIG_SIZE = 4 * 1024 + 256 + 128 + 512;
 
 constexpr static std::uint32_t IDLE_ERISC_L1_KERNEL_CONFIG_BASE = 32 * 1024;
 
-// config for 32 L1 buffers is at addr BUFFER_CONFIG_BASE
-// 12 bytes for each buffer: (addr, size, size_in_tiles)
-// addr and size are in 16B words (byte address >> 4)
-// this is a total of 32 * 3 * 4 = 384B
-constexpr static std::uint32_t CIRCULAR_BUFFER_CONFIG_BASE = L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE;
 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);
 
 constexpr static std::uint32_t PROFILER_L1_CONTROL_VECTOR_SIZE = 32;
 constexpr static std::uint32_t PROFILER_L1_CONTROL_BUFFER_SIZE = PROFILER_L1_CONTROL_VECTOR_SIZE * sizeof(uint32_t);
-constexpr static std::uint32_t PROFILER_L1_BUFFER_CONTROL = CIRCULAR_BUFFER_CONFIG_BASE + CIRCULAR_BUFFER_CONFIG_SIZE;
+constexpr static std::uint32_t PROFILER_L1_BUFFER_CONTROL = L1_KERNEL_CONFIG_BASE + L1_KERNEL_CONFIG_SIZE;
 
-// 4 uint32_t semaphores per core aligned to 16B
-constexpr static std::uint32_t SEMAPHORE_BASE = PROFILER_L1_BUFFER_CONTROL + PROFILER_L1_CONTROL_BUFFER_SIZE;
-constexpr static std::uint32_t NUM_SEMAPHORES = 8;
-constexpr static std::uint32_t SEMAPHORE_SIZE = NUM_SEMAPHORES * L1_ALIGNMENT;
+constexpr static std::uint32_t L1_UNRESERVED_BASE = ((PROFILER_L1_BUFFER_CONTROL + PROFILER_L1_CONTROL_BUFFER_SIZE - 1) | (DRAM_ALIGNMENT - 1)) + 1;
 
-constexpr static std::uint32_t L1_UNRESERVED_BASE = ((SEMAPHORE_BASE + SEMAPHORE_SIZE - 1) | (DRAM_ALIGNMENT - 1)) + 1;
 constexpr static std::uint32_t ERISC_L1_UNRESERVED_BASE = L1_UNRESERVED_BASE; // Start of unreserved space
 
 // Helper functions to convert NoC coordinates to NoC-0 coordinates, used in metal as "physical" coordinates.

tt_metal/hw/CMakeLists.txt

@@ -25,7 +25,7 @@ foreach(CORE IN LISTS CORES)
     add_custom_command(
         OUTPUT ${HW_OUTPUT_FILE}
         COMMAND ${CMAKE_COMMAND} -E make_directory ${HW_OUTPUT_DIR}
-        COMMAND ${CMAKE_CXX_COMPILER} -DLD_TARGET=${CORE_DEFINE} -DLD_TYPE=FIRMWARE -I${HW_INCLUDES} -E -P -x c -o ${HW_OUTPUT_FILE} ${CMAKE_CURRENT_SOURCE_DIR}/toolchain/main.ld
+        COMMAND ${CMAKE_CXX_COMPILER} -DLD_TARGET=${CORE_DEFINE} -DLD_TYPE=FIRMWARE -DCOMPILE_FOR_${CORE_DEFINE} -I${HW_INCLUDES} -E -P -x c -o ${HW_OUTPUT_FILE} ${CMAKE_CURRENT_SOURCE_DIR}/toolchain/main.ld
         DEPENDS ${CMAKE_CURRENT_SOURCE_DIR}/toolchain/main.ld ${CMAKE_CURRENT_SOURCE_DIR}/toolchain/memory.ld ${CMAKE_CURRENT_SOURCE_DIR}/toolchain/sections.ld ${DEV_MEM_MAP}
         COMMENT "Preprocessing toolchain/${CORE}.ld"
         VERBATIM