From 18462bef930afda699e4111b8bb280bad6113628 Mon Sep 17 00:00:00 2001 From: Paul Keller Date: Tue, 5 Dec 2023 22:09:00 +0000 Subject: [PATCH 1/4] #0: Reduce spew when logging not done cores --- tt_metal/llrt/llrt.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/tt_metal/llrt/llrt.cpp b/tt_metal/llrt/llrt.cpp index 1db5baa997a..33f5fab2de9 100644 --- a/tt_metal/llrt/llrt.cpp +++ b/tt_metal/llrt/llrt.cpp @@ -336,10 +336,10 @@ void wait_until_cores_done(chip_id_t device_id, std::unordered_set& not_done_phys_cores) { // poll the cores until the set of not done cores is empty - int loop_count = 0; + int loop_count = 1; while (!not_done_phys_cores.empty()) { // Print not-done cores - if (loop_count % 20 == 0) { + if (loop_count % 1000 == 0) { string not_done_cores_str = "Not done phys cores: "; for (const auto &core : not_done_phys_cores) { not_done_cores_str += (core.str() + " "); From e57c66c36ddd107f835c6608cbb6365d1b186056 Mon Sep 17 00:00:00 2001 From: Paul Keller Date: Tue, 5 Dec 2023 23:54:08 +0000 Subject: [PATCH 2/4] #4197: Improve thorw/assert/fatal message formatting Use log_fatal (reintroduced) to print both "fatal" and arguments of failure strings. --- tt_metal/common/assert.hpp | 28 +++++++++++++++++-- tt_metal/common/bfloat8.hpp | 2 +- tt_metal/common/logger.hpp | 10 +++++++ tt_metal/common/metal_soc_descriptor.cpp | 3 +- .../impl/buffers/circular_buffer_types.hpp | 10 +------ tt_metal/llrt/tt_cluster.cpp | 2 +- 6 files changed, 40 insertions(+), 15 deletions(-) diff --git a/tt_metal/common/assert.hpp b/tt_metal/common/assert.hpp index d9205b7dcfa..42844ef9bdd 100644 --- a/tt_metal/common/assert.hpp +++ b/tt_metal/common/assert.hpp @@ -104,6 +104,25 @@ void tt_assert_message(std::ostream& os, T const& t, Ts const&... ts) { tt_assert_message(os, ts...); } +template +void tt_assert_log_message(Ts const&... ts) { + std::string fmt; + for (int i = 0; i < sizeof...(ts); i++) { + fmt += "{} "; + } + log_fatal(fmt.c_str(), ts...); +} + +template +void tt_assert_log_message(const char *t, Ts const&... ts) { + log_fatal(t, ts...); +} + +template +void tt_assert_log_message(const std::string& t, Ts const&... ts) { + log_fatal(t.c_str(), ts...); +} + template [[ noreturn ]] void tt_throw(char const* file, int line, const std::string& assert_type, char const* condition_str, Ts const&... messages) { std::stringstream trace_message_ss = {}; @@ -111,13 +130,18 @@ template if constexpr (sizeof...(messages) > 0) { trace_message_ss << "info:" << std::endl; tt_assert_message(trace_message_ss, messages...); + tt_assert_log_message(messages...); } trace_message_ss << "backtrace:\n"; trace_message_ss << tt::assert::backtrace_to_string(100, 3, " --- "); trace_message_ss << std::flush; Logger::get().flush(); throw std::runtime_error(trace_message_ss.str()); +} +template +void tt_assert(char const* file, int line, const std::string& assert_type, char const* condition_str, Ts const&... messages) { + ::tt::assert::tt_throw(file, line, assert_type, condition_str, messages...); } template @@ -133,14 +157,14 @@ void tt_assert(char const* file, int line, const std::string& assert_type, bool // https://stackoverflow.com/questions/55933541/else-without-previous-if-error-when-defining-macro-with-arguments/55933720#55933720 #ifdef DEBUG #ifndef TT_ASSERT -#define TT_ASSERT(condition, ...) do{ if (not (condition)) ::tt::assert::tt_assert(__FILE__, __LINE__, "TT_ASSERT", (condition), #condition, ##__VA_ARGS__); }while(0) +#define TT_ASSERT(condition, ...) do{ if (not (condition)) tt::assert::tt_assert(__FILE__, __LINE__, "TT_ASSERT", (condition), #condition, ##__VA_ARGS__); } while(0) #endif #else #define TT_ASSERT(condition, ...) #endif #ifndef TT_THROW -#define TT_THROW(...) ::tt::assert::tt_throw(__FILE__, __LINE__, "TT_THROW", "tt::exception", ##__VA_ARGS__) +#define TT_THROW(...) tt::assert::tt_throw(__FILE__, __LINE__, "TT_THROW", "tt::exception", ##__VA_ARGS__) #endif #ifndef TT_FATAL diff --git a/tt_metal/common/bfloat8.hpp b/tt_metal/common/bfloat8.hpp index 01d42aba323..da1bdbdfbae 100644 --- a/tt_metal/common/bfloat8.hpp +++ b/tt_metal/common/bfloat8.hpp @@ -287,7 +287,7 @@ inline std::vector unpack_bfp8_tiles_into_float_vec(const std::vector(args)...); } +template +static void log_fatal(LogType type, char const* fmt, Args&&... args) { + Logger::get().log_level_type(Logger::Level::Fatal, type, fmt, std::forward(args)...); +} + +template +static void log_fatal(char const* fmt, Args&&... args) { + log_fatal(LogAlways, fmt, std::forward(args)...); +} + #undef LOGGER_TYPES } // namespace tt diff --git a/tt_metal/common/metal_soc_descriptor.cpp b/tt_metal/common/metal_soc_descriptor.cpp index 44bf3953731..3b762ffb42b 100644 --- a/tt_metal/common/metal_soc_descriptor.cpp +++ b/tt_metal/common/metal_soc_descriptor.cpp @@ -259,8 +259,7 @@ void metal_SocDescriptor::load_dispatch_and_banking_config(uint32_t harvesting_m uint32_t num_harvested_noc_rows = mask_bitset.count(); if (num_harvested_noc_rows > 2) { - TT_THROW( - tt::LogDevice, "At most two rows can be harvested, but detected {} harvested rows", num_harvested_noc_rows); + TT_THROW("At most two rows can be harvested, but detected {} harvested rows", num_harvested_noc_rows); } if (num_harvested_noc_rows == 1 and this->arch == tt::ARCH::GRAYSKULL) { TT_THROW("One row harvested Grayskull is not supported"); diff --git a/tt_metal/impl/buffers/circular_buffer_types.hpp b/tt_metal/impl/buffers/circular_buffer_types.hpp index 8bd384c0231..d0f03146d2f 100644 --- a/tt_metal/impl/buffers/circular_buffer_types.hpp +++ b/tt_metal/impl/buffers/circular_buffer_types.hpp @@ -39,7 +39,6 @@ class CircularBufferConfig { } if (total_size > buffer.size()) { TT_THROW( - tt::LogMetal, "Requested {} B but dynamic circular buffer cannot be larger than allocated L1 buffer of {} B", total_size, buffer.size()); @@ -50,28 +49,24 @@ class CircularBufferConfig { CircularBufferConfig set_page_size(uint8_t buffer_index, uint32_t page_size) { if (buffer_index > NUM_CIRCULAR_BUFFERS - 1) { TT_THROW( - tt::LogMetal, "Buffer index ({}) exceeds max number of circular buffers per core ({})", buffer_index, NUM_CIRCULAR_BUFFERS); } if (this->buffer_indices_.find(buffer_index) == this->buffer_indices_.end()) { TT_THROW( - tt::LogMetal, "Illegal circular buffer index {}. Page size can only be specified for buffer indices configured " "during config creation", buffer_index); } if (this->total_size_ % page_size != 0) { TT_THROW( - tt::LogMetal, "Total circular buffer size {} B must be divisible by page size {} B", this->total_size_, page_size); } if (page_size % sizeof(uint32_t) != 0) { - TT_THROW( - tt::LogMetal, "Page size must be divisible by sizeof(uint32_t) because buffers holds uint32_t values"); + TT_THROW("Page size must be divisible by sizeof(uint32_t) because buffers holds uint32_t values"); } this->page_sizes_[buffer_index] = page_size; @@ -81,7 +76,6 @@ class CircularBufferConfig { CircularBufferConfig set_total_size(uint32_t total_size) { if (dynamic_cb_ and total_size > this->max_size_.value()) { TT_THROW( - tt::LogMetal, "Cannot grow circular buffer to {} B. This is larger than associated dynamically allocated L1 buffer " "of {} B", total_size, @@ -118,7 +112,6 @@ class CircularBufferConfig { void set_config(const std::map &data_format_spec) { if (data_format_spec.size() > NUM_CIRCULAR_BUFFERS) { TT_THROW( - tt::LogMetal, "Only {} circular buffer slots are available but data formats are specified for {} indices", NUM_CIRCULAR_BUFFERS, data_format_spec.size()); @@ -127,7 +120,6 @@ class CircularBufferConfig { for (const auto &[buffer_index, data_format] : data_format_spec) { if (buffer_index > NUM_CIRCULAR_BUFFERS - 1) { TT_THROW( - tt::LogMetal, "Buffer index ({}) exceeds max number of circular buffers per core ({})", buffer_index, NUM_CIRCULAR_BUFFERS); diff --git a/tt_metal/llrt/tt_cluster.cpp b/tt_metal/llrt/tt_cluster.cpp index d97d1d74ea7..0773a8adcd0 100644 --- a/tt_metal/llrt/tt_cluster.cpp +++ b/tt_metal/llrt/tt_cluster.cpp @@ -309,7 +309,7 @@ tt_device &Cluster::get_driver(chip_id_t device_id) const { const metal_SocDescriptor &Cluster::get_soc_desc(chip_id_t chip) const { if (this->sdesc_per_chip_.find(chip) == this->sdesc_per_chip_.end()) { - TT_FATAL("Cannot access soc descriptor for {} before device driver is initialized! Call initialize_device_driver({}) first", chip, chip); + TT_THROW("Cannot access soc descriptor for {} before device driver is initialized! Call initialize_device_driver({}) first", chip, chip); } return this->sdesc_per_chip_.at(chip); } From f708f89dcaea8969387425c6cd42c5a94b4b0bdb Mon Sep 17 00:00:00 2001 From: Paul Keller Date: Wed, 6 Dec 2023 15:06:09 +0000 Subject: [PATCH 3/4] #4197: Make watcher error messages less terse --- tt_metal/llrt/watcher.cpp | 46 +++++++++++++++++++-------------------- 1 file changed, 23 insertions(+), 23 deletions(-) diff --git a/tt_metal/llrt/watcher.cpp b/tt_metal/llrt/watcher.cpp index 989cc63d2ae..16cbd9d1bc3 100644 --- a/tt_metal/llrt/watcher.cpp +++ b/tt_metal/llrt/watcher.cpp @@ -139,7 +139,7 @@ static const char * get_sanity_riscv_name(CoreCoord core, const launch_msg_t *la return "trisc2"; default: log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected riscv type on core {}: {}", core.str(), type); + TT_THROW("Watcher data corrupted, unexpected riscv type on core {}: {}", core.str(), type); } return nullptr; } @@ -156,7 +156,7 @@ static string get_debug_status(CoreCoord core, const launch_msg_t *launch_msg, c out += v; } else { log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected debug status on core {}, unprintable character {}", + TT_THROW("Watcher data corrupted, unexpected debug status on core {}, unprintable character {}", core.str(), (int)v); } } @@ -196,19 +196,19 @@ static void dump_noc_sanity_status(FILE *f, } break; case DebugSanitizeNocInvalidL1: - fprintf(f, "noc%d:%s{0x%08lx, %d}", noc, get_sanity_riscv_name(core, launch_msg, san->which), san->addr, san->len); + fprintf(f, "%s using noc%d reading L1[addr=0x%08lx,len=%d]\n", get_sanity_riscv_name(core, launch_msg, san->which), noc, san->addr, san->len); fflush(f); log_running_kernels(launch_msg); log_info("Watcher stopped the device due to bad NOC L1/reg address"); log_waypoint(core, launch_msg, debug_status); - snprintf(buf, sizeof(buf), "On core %s: noc%d:%s{0x%08lx, %d}", - core.str().c_str(), noc, get_sanity_riscv_name(core, launch_msg, san->which), san->addr, san->len); + snprintf(buf, sizeof(buf), "On core %s: %s using noc%d reading L1[addr=0x%08lx,len=%d]", + core.str().c_str(), get_sanity_riscv_name(core, launch_msg, san->which), noc, san->addr, san->len); TT_THROW(buf); break; case DebugSanitizeNocInvalidUnicast: - fprintf(f, "noc%d:%s{(%02ld,%02ld) 0x%08lx, %d}", - noc, + fprintf(f, "%s using noc%d tried to access core (%02ld,%02ld) L1[addr=0x%08lx,len=%d]\n", get_sanity_riscv_name(core, launch_msg, san->which), + noc, NOC_UNICAST_ADDR_X(san->addr), NOC_UNICAST_ADDR_Y(san->addr), NOC_LOCAL_ADDR_OFFSET(san->addr), san->len); @@ -216,19 +216,19 @@ static void dump_noc_sanity_status(FILE *f, log_info("Watcher stopped the device due to bad NOC unicast transaction"); log_running_kernels(launch_msg); log_waypoint(core, launch_msg, debug_status); - snprintf(buf, sizeof(buf), "On core %s: noc%d:%s{(%02ld,%02ld) 0x%08lx, %d}", + snprintf(buf, sizeof(buf), "On core %s: %s using noc%d tried to accesss core (%02ld,%02ld) L1[addr=0x%08lx,len=%d]", core.str().c_str(), - noc, get_sanity_riscv_name(core, launch_msg, san->which), + noc, NOC_UNICAST_ADDR_X(san->addr), NOC_UNICAST_ADDR_Y(san->addr), NOC_LOCAL_ADDR_OFFSET(san->addr), san->len); TT_THROW(buf); break; case DebugSanitizeNocInvalidMulticast: - fprintf(f, "noc%d:%s{(%02ld,%02ld)-(%02ld,%02ld) 0x%08lx, %d}", - noc, + fprintf(f, "%s using noc%d tried to access core range (%02ld,%02ld)-(%02ld,%02ld) L1[addr=0x%08lx,len=%d]\n", get_sanity_riscv_name(core, launch_msg, san->which), + noc, NOC_MCAST_ADDR_START_X(san->addr), NOC_MCAST_ADDR_START_Y(san->addr), NOC_MCAST_ADDR_END_X(san->addr), @@ -238,10 +238,10 @@ static void dump_noc_sanity_status(FILE *f, log_info("Watcher stopped the device due to bad NOC multicast transaction"); log_running_kernels(launch_msg); log_waypoint(core, launch_msg, debug_status); - snprintf(buf, sizeof(buf), "On core %s: noc%d:%s{(%02ld,%02ld)-(%02ld,%02ld) 0x%08lx, %d}", + snprintf(buf, sizeof(buf), "On core %s: %s using noc%d tried to access core range (%02ld,%02ld)-(%02ld,%02ld) L1[addr=0x%08lx,len=%d]}", core.str().c_str(), - noc, get_sanity_riscv_name(core, launch_msg, san->which), + noc, NOC_MCAST_ADDR_START_X(san->addr), NOC_MCAST_ADDR_START_Y(san->addr), NOC_MCAST_ADDR_END_X(san->addr), @@ -251,7 +251,7 @@ static void dump_noc_sanity_status(FILE *f, break; default: log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected noc debug state on core {}, unknown failure code: {}\n", + TT_THROW("Watcher unexpected data corruption, noc debug state on core {}, unknown failure code: {}\n", core.str(), san->invalid); } } @@ -274,7 +274,7 @@ static void dump_run_state(FILE *f, CoreCoord core, const launch_msg_t *launch_m else if (state == RUN_MSG_DONE) code = 'D'; if (code == 'U') { log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected run state on core{}: {} (expected {} or {} or {})", + TT_THROW("Watcher data corruption, unexpected run state on core{}: {} (expected {} or {} or {})", core.str(), state, RUN_MSG_INIT, RUN_MSG_GO, RUN_MSG_DONE); } else { fprintf(f, "%c", code); @@ -294,7 +294,7 @@ static void dump_run_mailboxes(FILE *f, fprintf(f, "H"); } else { log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected launch mode on core {}: {} (expected {} or {})", + TT_THROW("Watcher data corruption, unexpected launch mode on core {}: {} (expected {} or {})", core.str(), launch_msg->mode, DISPATCH_MODE_DEV, DISPATCH_MODE_HOST); } @@ -302,7 +302,7 @@ static void dump_run_mailboxes(FILE *f, fprintf(f, "%d", launch_msg->brisc_noc_id); } else { log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected brisc noc_id on core {}: {} (expected 0 or 1)", + TT_THROW("Watcher data corruption, unexpected brisc noc_id on core {}: {} (expected 0 or 1)", core.str(), launch_msg->brisc_noc_id); } @@ -316,7 +316,7 @@ static void dump_run_mailboxes(FILE *f, fprintf(f, "b"); } else { log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected brisc enable on core {}: {} (expected 0 or 1)", + TT_THROW("Watcher data corruption, unexpected brisc enable on core {}: {} (expected 0 or 1)", core.str(), launch_msg->enable_brisc); } @@ -327,7 +327,7 @@ static void dump_run_mailboxes(FILE *f, fprintf(f, "n"); } else { log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected ncrisc enable on core {}: {} (expected 0 or 1)", + TT_THROW("Watcher data corruption, unexpected ncrisc enable on core {}: {} (expected 0 or 1)", core.str(), launch_msg->enable_ncrisc); } @@ -338,7 +338,7 @@ static void dump_run_mailboxes(FILE *f, fprintf(f, "t"); } else { log_running_kernels(launch_msg); - TT_THROW("Watcher unexpected trisc enable on core {}: {} (expected 0 or 1)", + TT_THROW("Watcher data corruption, unexpected trisc enable on core {}: {} (expected 0 or 1)", core.str(), launch_msg->enable_triscs); } @@ -389,19 +389,19 @@ static void validate_kernel_ids(FILE *f, const launch_msg_t *launch) { if (launch->brisc_watcher_kernel_id >= kernel_names.size()) { - TT_THROW("Watcher unexpected brisc kernel id on core {}: {} (last valid {})", + TT_THROW("Watcher data corruption, unexpected brisc kernel id on core {}: {} (last valid {})", core.str(), launch->brisc_watcher_kernel_id, kernel_names.size()); } used_kernel_names[launch->brisc_watcher_kernel_id] = true; if (launch->ncrisc_watcher_kernel_id >= kernel_names.size()) { - TT_THROW("Watcher unexpected ncrisc kernel id on core {}: {} (last valid {})", + TT_THROW("Watcher data corruption, unexpected ncrisc kernel id on core {}: {} (last valid {})", core.str(), launch->ncrisc_watcher_kernel_id, kernel_names.size()); } used_kernel_names[launch->ncrisc_watcher_kernel_id] = true; if (launch->triscs_watcher_kernel_id >= kernel_names.size()) { - TT_THROW("Watcher unexpected trisc kernel id on core {}: {} (last valid {})", + TT_THROW("Watcher data corruption, unexpected trisc kernel id on core {}: {} (last valid {})", core.str(), launch->triscs_watcher_kernel_id, kernel_names.size()); } used_kernel_names[launch->triscs_watcher_kernel_id] = true; From 1bcc5734a9f62cfaac9eba5e2654e244aaa7b898 Mon Sep 17 00:00:00 2001 From: Paul Keller Date: Wed, 6 Dec 2023 18:26:44 +0000 Subject: [PATCH 4/4] #4197: Fix tests that use assert instead of throw Assert can go away w/ release builds leading to unused var warnings. --- .../tt_dnn/op_library/moreh_layernorm/moreh_layernorm_op.cpp | 2 +- .../moreh_layernorm_backward_gamma_beta_grad.cpp | 2 +- .../input_grad/moreh_layernorm_backward_input_grad.cpp | 2 +- .../moreh_matmul/multi_core/moreh_matmul_op_multi_core.cpp | 2 +- .../moreh_matmul_backward/sum/moreh_sum_multi_core.cpp | 2 +- .../moreh_softmax/softmax_c_large/softmax_c_large.cpp | 2 +- .../moreh_softmax/softmax_h_large/softmax_h_large.cpp | 2 +- .../moreh_softmax/softmax_h_small/softmax_h_small.cpp | 2 +- .../moreh_softmax/softmax_w_large/softmax_w_large.cpp | 2 +- .../moreh_softmax/softmax_w_small/softmax_w_small.cpp | 2 +- .../softmax_backward_c_large/softmax_backward_c_large.cpp | 2 +- .../softmax_backward_h_large/softmax_backward_h_large.cpp | 2 +- .../softmax_backward_h_small/softmax_backward_h_small.cpp | 2 +- .../softmax_backward_w_large/softmax_backward_w_large.cpp | 2 +- .../softmax_backward_w_small/softmax_backward_w_small.cpp | 2 +- 15 files changed, 15 insertions(+), 15 deletions(-) diff --git a/tt_eager/tt_dnn/op_library/moreh_layernorm/moreh_layernorm_op.cpp b/tt_eager/tt_dnn/op_library/moreh_layernorm/moreh_layernorm_op.cpp index 9c28adb70b6..d48c90cdb05 100644 --- a/tt_eager/tt_dnn/op_library/moreh_layernorm/moreh_layernorm_op.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_layernorm/moreh_layernorm_op.cpp @@ -347,7 +347,7 @@ operation::ProgramWithCallbacks moreh_layernorm_impl( } else if (core_group_2.core_coord_in_core_ranges(core)) { num_rows_per_core = num_rows_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges."); + TT_THROW("Core not in specified core ranges."); } const std::vector reader_runtime_args{ diff --git a/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/gamma_beta_grad/moreh_layernorm_backward_gamma_beta_grad.cpp b/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/gamma_beta_grad/moreh_layernorm_backward_gamma_beta_grad.cpp index ff6a986eebb..fc76c5cdb91 100644 --- a/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/gamma_beta_grad/moreh_layernorm_backward_gamma_beta_grad.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/gamma_beta_grad/moreh_layernorm_backward_gamma_beta_grad.cpp @@ -236,7 +236,7 @@ operation::ProgramWithCallbacks moreh_layernorm_backward_gamma_beta_grad_impl( } else if (core_group_2.core_coord_in_core_ranges(core)) { num_cols_per_core = num_cols_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges."); + TT_THROW("Core not in specified core ranges."); } const std::vector reader_runtime_args{ diff --git a/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/input_grad/moreh_layernorm_backward_input_grad.cpp b/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/input_grad/moreh_layernorm_backward_input_grad.cpp index 8a07393ea74..b26400b30b0 100644 --- a/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/input_grad/moreh_layernorm_backward_input_grad.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_layernorm_backward/input_grad/moreh_layernorm_backward_input_grad.cpp @@ -277,7 +277,7 @@ operation::ProgramWithCallbacks moreh_layernorm_backward_input_grad_impl( } else if (core_group_2.core_coord_in_core_ranges(core)) { num_rows_per_core = num_rows_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges."); + TT_THROW("Core not in specified core ranges."); } const std::vector reader_runtime_args{ diff --git a/tt_eager/tt_dnn/op_library/moreh_matmul/multi_core/moreh_matmul_op_multi_core.cpp b/tt_eager/tt_dnn/op_library/moreh_matmul/multi_core/moreh_matmul_op_multi_core.cpp index 0bf1d1eae73..abb7044dadc 100644 --- a/tt_eager/tt_dnn/op_library/moreh_matmul/multi_core/moreh_matmul_op_multi_core.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_matmul/multi_core/moreh_matmul_op_multi_core.cpp @@ -184,7 +184,7 @@ operation::ProgramWithCallbacks moreh_matmul_multi_core( } else if (core_group_2.core_coord_in_core_ranges(core)) { num_output_tiles_per_core = num_output_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } tt_metal::SetRuntimeArgs( diff --git a/tt_eager/tt_dnn/op_library/moreh_matmul_backward/sum/moreh_sum_multi_core.cpp b/tt_eager/tt_dnn/op_library/moreh_matmul_backward/sum/moreh_sum_multi_core.cpp index 80d00563fbd..cea0a9eee83 100644 --- a/tt_eager/tt_dnn/op_library/moreh_matmul_backward/sum/moreh_sum_multi_core.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_matmul_backward/sum/moreh_sum_multi_core.cpp @@ -130,7 +130,7 @@ operation::ProgramWithCallbacks moreh_sum_multi_core(const Tensor &src, const Te } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_cols_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges."); + TT_THROW("Core not in specified core ranges."); } SetRuntimeArgs( diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_c_large/softmax_c_large.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_c_large/softmax_c_large.cpp index a12acc72606..34467bab867 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_c_large/softmax_c_large.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_c_large/softmax_c_large.cpp @@ -107,7 +107,7 @@ operation::ProgramWithCallbacks moreh_softmax_c_large(const Tensor &input, Tenso } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } vector reader_args = { diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_large/softmax_h_large.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_large/softmax_h_large.cpp index c6c91d91fc2..45adf63f7f5 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_large/softmax_h_large.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_large/softmax_h_large.cpp @@ -93,7 +93,7 @@ operation::ProgramWithCallbacks moreh_softmax_h_large(const Tensor &input, Tenso } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f; diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_small/softmax_h_small.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_small/softmax_h_small.cpp index 914b9c2b3b0..922a9d09b37 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_small/softmax_h_small.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_h_small/softmax_h_small.cpp @@ -115,7 +115,7 @@ operation::ProgramWithCallbacks moreh_softmax_h_small(const Tensor &input, Tenso } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f; diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_large/softmax_w_large.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_large/softmax_w_large.cpp index b39139ece1d..b4febad413a 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_large/softmax_w_large.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_large/softmax_w_large.cpp @@ -93,7 +93,7 @@ operation::ProgramWithCallbacks moreh_softmax_w_large(const Tensor &input, Tenso } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f; diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_small/softmax_w_small.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_small/softmax_w_small.cpp index b3deb62fcb4..e775b0d9430 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_small/softmax_w_small.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax/softmax_w_small/softmax_w_small.cpp @@ -114,7 +114,7 @@ operation::ProgramWithCallbacks moreh_softmax_w_small(const Tensor &input, Tenso } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f; diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_c_large/softmax_backward_c_large.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_c_large/softmax_backward_c_large.cpp index ef525110dc4..64e1ba842eb 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_c_large/softmax_backward_c_large.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_c_large/softmax_backward_c_large.cpp @@ -110,7 +110,7 @@ operation::ProgramWithCallbacks moreh_softmax_backward_c_large(const Tensor &out } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } vector reader_args = { diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_large/softmax_backward_h_large.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_large/softmax_backward_h_large.cpp index 11a398a0847..f59cc4ed005 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_large/softmax_backward_h_large.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_large/softmax_backward_h_large.cpp @@ -97,7 +97,7 @@ operation::ProgramWithCallbacks moreh_softmax_backward_h_large(const Tensor &out } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f; diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_small/softmax_backward_h_small.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_small/softmax_backward_h_small.cpp index 33b331c58d9..af1ef2106e5 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_small/softmax_backward_h_small.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_h_small/softmax_backward_h_small.cpp @@ -120,7 +120,7 @@ operation::ProgramWithCallbacks moreh_softmax_backward_h_small(const Tensor &out } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f; diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_large/softmax_backward_w_large.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_large/softmax_backward_w_large.cpp index dc60b4b4fc0..05ddb4a0dc4 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_large/softmax_backward_w_large.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_large/softmax_backward_w_large.cpp @@ -96,7 +96,7 @@ operation::ProgramWithCallbacks moreh_softmax_backward_w_large(const Tensor &out } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f; diff --git a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_small/softmax_backward_w_small.cpp b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_small/softmax_backward_w_small.cpp index 0b45aa86e71..00cd056655c 100644 --- a/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_small/softmax_backward_w_small.cpp +++ b/tt_eager/tt_dnn/op_library/moreh_softmax_backward/softmax_backward_w_small/softmax_backward_w_small.cpp @@ -120,7 +120,7 @@ operation::ProgramWithCallbacks moreh_softmax_backward_w_small(const Tensor &out } else if (core_group_2.core_coord_in_core_ranges(core)) { num_tiles_per_core = num_tiles_per_core_group_2; } else { - TT_ASSERT(false, "Core not in specified core ranges"); + TT_THROW("Core not in specified core ranges"); } float scaler = 1.0f;