Skip to content

Commit

Permalink
Move NOC_0_X/Y behind Hal
Browse files Browse the repository at this point in the history
  • Loading branch information
@blozano-tt
blozano-tt committed Nov 9, 2024
1 parent b27fb15 commit 4f2fff5
Show file tree
Hide file tree
Showing 7 changed files with 43 additions and 58 deletions.
18 changes: 9 additions & 9 deletions tests/tt_metal/tt_metal/perf_microbenchmark/dispatch/common.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,7 @@
#include "tt_metal/impl/dispatch/cq_commands.hpp"
#include "noc/noc_parameters.h"

#include "tt_metal/hostdevcommon/common_runtime_address_map.h" // NOC_0_X
#include "tt_metal/llrt/hal.hpp"

extern bool debug_g;
extern bool use_coherent_data_g;
Expand Down Expand Up @@ -490,15 +490,15 @@ void configure_kernel_variant(
const auto& grid_size = device->grid_size();

std::map<string, string> defines = {
{"MY_NOC_X", std::to_string(NOC_0_X(my_noc_index, grid_size.x, phys_my_core.x))},
{"MY_NOC_Y", std::to_string(NOC_0_Y(my_noc_index, grid_size.y, phys_my_core.y))},
{"MY_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.x, phys_my_core.x))},
{"MY_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.y, phys_my_core.y))},
{"UPSTREAM_NOC_INDEX", std::to_string(upstream_noc_index)},
{"UPSTREAM_NOC_X", std::to_string(NOC_0_X(upstream_noc_index, grid_size.x, phys_upstream_core.x))},
{"UPSTREAM_NOC_Y", std::to_string(NOC_0_Y(upstream_noc_index, grid_size.y, phys_upstream_core.y))},
{"DOWNSTREAM_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, phys_downstream_core.x))},
{"DOWNSTREAM_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, phys_downstream_core.y))},
{"DOWNSTREAM_SLAVE_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, 0xff))},
{"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, 0xff))}, // todo, add testing with dispatch_s once it processes more than go signals
{"UPSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.x, phys_upstream_core.x))},
{"UPSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.y, phys_upstream_core.y))},
{"DOWNSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, phys_downstream_core.x))},
{"DOWNSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, phys_downstream_core.y))},
{"DOWNSTREAM_SLAVE_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, 0xff))},
{"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, 0xff))}, // todo, add testing with dispatch_s once it processes more than go signals
{"FD_CORE_TYPE", std::to_string(0)}, // todo, support dispatch on eth
};
compile_args.push_back(is_dram_variant);
Expand Down
9 changes: 6 additions & 3 deletions tests/tt_metal/tt_metal/unit_tests_common/watcher/test_noc_sanitize.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8,7 +8,10 @@
#include "tt_metal/detail/tt_metal.hpp"
#include "tt_metal/host_api.hpp"
#include "common/bfloat16.hpp"
#include "hostdevcommon/common_runtime_address_map.h"

// Do we really want to expose Hal like this?
// This looks like an API level test
#include "llrt/hal.hpp"

//////////////////////////////////////////////////////////////////////////////////////////
// A test for checking watcher NOC sanitization.
Expand Down Expand Up @@ -155,8 +158,8 @@ void RunTestOnCore(WatcherFixture* fixture, Device* device, CoreCoord &core, boo
const metal_SocDescriptor& soc_d = tt::Cluster::instance().get_soc_desc(device->id());
int noc = (use_ncrisc) ? 1 : 0;
CoreCoord target_phys_core = {
NOC_0_X(noc, soc_d.grid_size.x, input_dram_noc_xy.x),
NOC_0_Y(noc, soc_d.grid_size.y, input_dram_noc_xy.y)
tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.x, input_dram_noc_xy.x),
tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.y, input_dram_noc_xy.y)
};
string risc_name = (is_eth_core) ? "erisc" : "brisc";
if (use_ncrisc)
Expand Down
4 changes: 0 additions & 4 deletions tt_metal/hostdevcommon/common_runtime_address_map.h
View file
Original file line number Diff line number Diff line change
Expand Up @@ -17,8 +17,4 @@
constexpr static std::uint32_t L1_KERNEL_CONFIG_BASE = MEM_MAP_END;
constexpr static std::uint32_t L1_KERNEL_CONFIG_SIZE = 69 * 1024;

// Helper functions to convert NoC coordinates to NoC-0 coordinates, used in metal as "physical" coordinates.
#define NOC_0_X(noc_index, noc_size_x, x) (noc_index == 0 ? (x) : (noc_size_x-1-(x)))
#define NOC_0_Y(noc_index, noc_size_y, y) (noc_index == 0 ? (y) : (noc_size_y-1-(y)))

static_assert(L1_KERNEL_CONFIG_BASE % L1_ALIGNMENT == 0);
3 changes: 1 addition & 2 deletions tt_metal/impl/debug/watcher_device_reader.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -16,7 +16,6 @@
// FIXME: Avoid dependence on ARCH_NAME specific includes
#include "dev_mem_map.h" // for MEM_BRISC_STAC...
#include "eth_l1_address_map.h" // for address_map
#include "hostdevcommon/common_runtime_address_map.h" // for NOC_0_X, NOC_0_Y
#include "hw/inc/dev_msgs.h"

#include "third_party/umd/device/tt_arch_types.h"
Expand Down Expand Up @@ -79,7 +78,7 @@ static string get_noc_target_str(Device *device, CoreDescriptor &core, int noc,
// Get the physical coord from the noc coord
const metal_SocDescriptor &soc_d = tt::Cluster::instance().get_soc_desc(device->id());
CoreCoord phys_core = {
NOC_0_X(noc, soc_d.grid_size.x, noc_coord.x), NOC_0_Y(noc, soc_d.grid_size.y, noc_coord.y)};
tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.x, noc_coord.x), tt::tt_metal::hal.noc_coordinate(noc, soc_d.grid_size.y, noc_coord.y)};

CoreType core_type;
try {
Expand Down
38 changes: 18 additions & 20 deletions tt_metal/impl/device/device.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -24,8 +24,6 @@
#include "tt_metal/tools/profiler/tt_metal_tracy.hpp"
#include "llrt/hal.hpp"

#include "tt_metal/hostdevcommon/common_runtime_address_map.h" // NOC_0_X

namespace tt {

namespace tt_metal {
Expand Down Expand Up @@ -769,15 +767,15 @@ void Device::configure_kernel_variant(

std::map<string, string> defines = {
{"DISPATCH_KERNEL", "1"},
{"MY_NOC_X", std::to_string(NOC_0_X(my_noc_index, grid_size.x, kernel_physical_core.x))},
{"MY_NOC_Y", std::to_string(NOC_0_Y(my_noc_index, grid_size.y, kernel_physical_core.y))},
{"MY_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.x, kernel_physical_core.x))},
{"MY_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(my_noc_index, grid_size.y, kernel_physical_core.y))},
{"UPSTREAM_NOC_INDEX", std::to_string(upstream_noc_index)},
{"UPSTREAM_NOC_X", std::to_string(NOC_0_X(upstream_noc_index, grid_size.x, upstream_physical_core.x))},
{"UPSTREAM_NOC_Y", std::to_string(NOC_0_Y(upstream_noc_index, grid_size.y, upstream_physical_core.y))},
{"DOWNSTREAM_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, downstream_physical_core.x))},
{"DOWNSTREAM_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, downstream_physical_core.y))},
{"DOWNSTREAM_SLAVE_NOC_X", std::to_string(NOC_0_X(downstream_noc_index, grid_size.x, downstream_slave_physical_core.x))},
{"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(NOC_0_Y(downstream_noc_index, grid_size.y, downstream_slave_physical_core.y))},
{"UPSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.x, upstream_physical_core.x))},
{"UPSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(upstream_noc_index, grid_size.y, upstream_physical_core.y))},
{"DOWNSTREAM_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, downstream_physical_core.x))},
{"DOWNSTREAM_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, downstream_physical_core.y))},
{"DOWNSTREAM_SLAVE_NOC_X", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.x, downstream_slave_physical_core.x))},
{"DOWNSTREAM_SLAVE_NOC_Y", std::to_string(tt::tt_metal::hal.noc_coordinate(downstream_noc_index, grid_size.y, downstream_slave_physical_core.y))},
{"FD_CORE_TYPE", std::to_string(programmable_core_type_index)},
};
if (force_watcher_no_inline) {
Expand Down Expand Up @@ -3099,8 +3097,8 @@ std::vector<CoreCoord> Device::ethernet_cores_from_logical_cores(const std::vect
uint32_t Device::get_noc_unicast_encoding(uint8_t noc_index, const CoreCoord& physical_core) const {
const auto& grid_size = this->grid_size();
return NOC_XY_ENCODING(
NOC_0_X(noc_index, grid_size.x, physical_core.x),
NOC_0_Y(noc_index, grid_size.y, physical_core.y)
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_core.x),
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_core.y)
);
}

Expand All @@ -3110,17 +3108,17 @@ uint32_t Device::get_noc_multicast_encoding(uint8_t noc_index, const CoreRange&
// NOC 1 mcasts from bottom left to top right, so we need to reverse the coords
if (noc_index == 0) {
return NOC_MULTICAST_ENCODING(
NOC_0_X(noc_index, grid_size.x, physical_cores.start_coord.x),
NOC_0_Y(noc_index, grid_size.y, physical_cores.start_coord.y),
NOC_0_X(noc_index, grid_size.x, physical_cores.end_coord.x),
NOC_0_Y(noc_index, grid_size.y, physical_cores.end_coord.y)
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.start_coord.x),
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.start_coord.y),
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.end_coord.x),
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.end_coord.y)
);
} else {
return NOC_MULTICAST_ENCODING(
NOC_0_X(noc_index, grid_size.x, physical_cores.end_coord.x),
NOC_0_Y(noc_index, grid_size.y, physical_cores.end_coord.y),
NOC_0_X(noc_index, grid_size.x, physical_cores.start_coord.x),
NOC_0_Y(noc_index, grid_size.y, physical_cores.start_coord.y)
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.end_coord.x),
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.end_coord.y),
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.x, physical_cores.start_coord.x),
tt::tt_metal::hal.noc_coordinate(noc_index, grid_size.y, physical_cores.start_coord.y)
);
}
}
Expand Down
24 changes: 4 additions & 20 deletions tt_metal/jit_build/genfiles.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -11,7 +11,6 @@

#include "common/tt_backend_api_types.hpp"
#include "common/utils.hpp"
#include "hostdevcommon/common_runtime_address_map.h" // NOC_0_X
#include "hostdevcommon/common_values.hpp"
#include "jit_build/build.hpp"
#include "jit_build/settings.hpp"
Expand Down Expand Up @@ -588,8 +587,8 @@ std::string generate_bank_to_noc_coord_descriptor_string(
ss << " {"
<< "\t// noc=" << noc << endl;
for (unsigned int bank_id = 0; bank_id < dram_bank_map.size(); bank_id++) {
uint16_t noc_x = NOC_0_X(noc, grid_size.x, dram_bank_map[bank_id].x);
uint16_t noc_y = NOC_0_Y(noc, grid_size.y, dram_bank_map[bank_id].y);
uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, grid_size.x, dram_bank_map[bank_id].x);
uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, grid_size.y, dram_bank_map[bank_id].y);
ss << " (((" << noc_y << " << NOC_ADDR_NODE_ID_BITS) | " << noc_x << ") << NOC_COORD_REG_OFFSET),"
<< "\t// NOC_X=" << noc_x << " NOC_Y=" << noc_y << endl;
}
Expand All @@ -609,8 +608,8 @@ std::string generate_bank_to_noc_coord_descriptor_string(
ss << " {"
<< "\t// noc=" << noc << endl;
for (unsigned int bank_id = 0; bank_id < l1_bank_map.size(); bank_id++) {
uint16_t noc_x = NOC_0_X(noc, grid_size.x, l1_bank_map[bank_id].x);
uint16_t noc_y = NOC_0_Y(noc, grid_size.y, l1_bank_map[bank_id].y);
uint16_t noc_x = tt::tt_metal::hal.noc_coordinate(noc, grid_size.x, l1_bank_map[bank_id].x);
uint16_t noc_y = tt::tt_metal::hal.noc_coordinate(noc, grid_size.y, l1_bank_map[bank_id].y);
ss << " (((" << noc_y << " << NOC_ADDR_NODE_ID_BITS) | " << noc_x << ") << NOC_COORD_REG_OFFSET),"
<< "\t// NOC_X=" << noc_x << " NOC_Y=" << noc_y << endl;
}
Expand Down Expand Up @@ -667,19 +666,4 @@ void jit_build_genfiles_bank_to_noc_coord_descriptor(
file_stream_siec.close();
}

static string generate_noc_core_xy_range_define(const std::vector<CoreCoord>& cores) {
stringstream ss;

string end_of_line = " \\\n ( \\";
for (const auto& core : cores) {
ss << end_of_line << endl;
ss << " ((x) == NOC_0_X(noc_idx, noc_size_x, (uint32_t)" << core.x
<< ") && (y) == NOC_0_Y(noc_idx, noc_size_y, (uint32_t)" << core.y << "))";
end_of_line = " || \\";
}
ss << ")" << endl;

return ss.str();
}

} // namespace tt::tt_metal
5 changes: 5 additions & 0 deletions tt_metal/llrt/hal.hpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -137,6 +137,11 @@ class Hal {

tt::ARCH get_arch() {return arch_;}

template <typename IndexType, typename SizeType, typename CoordType>
auto noc_coordinate(IndexType noc_index, SizeType noc_size, CoordType coord) const -> decltype(noc_size - 1 - coord) {
return noc_index == 0 ? coord : (noc_size - 1 - coord);
}

uint32_t get_programmable_core_type_count() const;
HalProgrammableCoreType get_programmable_core_type(uint32_t core_type_index) const;
uint32_t get_programmable_core_type_index(HalProgrammableCoreType programmable_core_type_index) const;
Expand Down

0 comments on commit 4f2fff5

Please sign in to comment.