diff --git a/include/ttmlir/Target/TTNN/program.fbs b/include/ttmlir/Target/TTNN/program.fbs index f145aaf65..19b1dbc92 100644 --- a/include/ttmlir/Target/TTNN/program.fbs +++ b/include/ttmlir/Target/TTNN/program.fbs @@ -286,6 +286,7 @@ union OpType { table Operation { type: OpType; debug_info: string; + loc_info: string; } table Program { diff --git a/include/ttmlir/Target/Utils/FuncOpToProgram.h b/include/ttmlir/Target/Utils/FuncOpToProgram.h index d9e8d9820..a28f2f5e9 100644 --- a/include/ttmlir/Target/Utils/FuncOpToProgram.h +++ b/include/ttmlir/Target/Utils/FuncOpToProgram.h @@ -31,6 +31,13 @@ inline std::string getOpDebugString(mlir::Operation *op, return str; }; +inline std::string getOpLocInfo(mlir::Operation *op) { + std::string str; + llvm::raw_string_ostream os(str); + op->getLoc().print(os); + return str; +} + inline Value getOperandThroughDPSOps(Value value) { auto *op = value.getDefiningOp(); if (!op) { @@ -76,7 +83,8 @@ Program funcOpToProgram(FlatbufferObjectCache &cache, func::FuncOp entry, } } else { std::string debugStr = getOpDebugString(op, printFlags); - program.ops.push_back(fn(cache, op, debugStr)); + std::string locInfo = getOpLocInfo(op); + program.ops.push_back(fn(cache, op, debugStr, locInfo)); } }); diff --git a/include/ttmlir/Target/Utils/MLIRToFlatbuffer.h b/include/ttmlir/Target/Utils/MLIRToFlatbuffer.h index d5be2bb97..cb9439d97 100644 --- a/include/ttmlir/Target/Utils/MLIRToFlatbuffer.h +++ b/include/ttmlir/Target/Utils/MLIRToFlatbuffer.h @@ -479,7 +479,11 @@ toDebugInfo(::flatbuffers::FlatBufferBuilder &fbb, std::string const &name, ModuleOp module) { std::string source; llvm::raw_string_ostream os(source); - module->print(os); + + mlir::OpPrintingFlags flags; + flags.enableDebugInfo(); // Enable the loc dumping + module->print(os, flags); + return ::tt::target::CreateMLIRDirect(fbb, name.c_str(), source.c_str()); } } // namespace mlir::tt diff --git a/lib/Target/TTNN/TTNNToFlatbuffer.cpp b/lib/Target/TTNN/TTNNToFlatbuffer.cpp index 34a0c4725..9706880e3 100644 --- a/lib/Target/TTNN/TTNNToFlatbuffer.cpp +++ b/lib/Target/TTNN/TTNNToFlatbuffer.cpp @@ -163,10 +163,10 @@ createDeviceRef(FlatbufferObjectCache &cache, Value device) { template ::flatbuffers::Offset<::tt::target::ttnn::Operation> createOperation(FlatbufferObjectCache &cache, ::flatbuffers::Offset op, - std::string const &debugString) { + std::string const &debugString, std::string const &locInfo) { return CreateOperationDirect( *cache.fbb, ::tt::target::ttnn::OpTypeTraits::enum_value, op.Union(), - debugString.c_str()); + debugString.c_str(), locInfo.c_str()); } ::flatbuffers::Offset<::tt::target::ttnn::GetDeviceOp> @@ -701,218 +701,263 @@ createDeallocateOp(FlatbufferObjectCache &cache, DeallocateOp op) { ::flatbuffers::Offset<::tt::target::ttnn::Operation> emitTTNNOperation(FlatbufferObjectCache &cache, Operation *op, - std::string const &debugString) { + std::string const &debugString, std::string const &locInfo) { if (auto getDeviceOp = dyn_cast(op); getDeviceOp) { - return createOperation(cache, createOp(cache, getDeviceOp), debugString); + return createOperation(cache, createOp(cache, getDeviceOp), debugString, + locInfo); } if (auto toMemoryConfigOp = dyn_cast(op); toMemoryConfigOp) { return createOperation(cache, createOp(cache, toMemoryConfigOp), - debugString); + debugString, locInfo); } if (auto toLayoutOp = dyn_cast(op); toLayoutOp) { - return createOperation(cache, createOp(cache, toLayoutOp), debugString); + return createOperation(cache, createOp(cache, toLayoutOp), debugString, + locInfo); } if (auto typecastOp = dyn_cast(op); typecastOp) { - return createOperation(cache, createOp(cache, typecastOp), debugString); + return createOperation(cache, createOp(cache, typecastOp), debugString, + locInfo); } if (auto toDeviceOp = dyn_cast(op); toDeviceOp) { - return createOperation(cache, createOp(cache, toDeviceOp), debugString); + return createOperation(cache, createOp(cache, toDeviceOp), debugString, + locInfo); } if (auto fromDeviceOp = dyn_cast(op); fromDeviceOp) { - return createOperation(cache, createOp(cache, fromDeviceOp), debugString); + return createOperation(cache, createOp(cache, fromDeviceOp), debugString, + locInfo); } if (auto emptyOp = dyn_cast(op); emptyOp) { - return createOperation(cache, createOp(cache, emptyOp), debugString); + return createOperation(cache, createOp(cache, emptyOp), debugString, + locInfo); } if (auto fullOp = dyn_cast(op); fullOp) { - return createOperation(cache, createOp(cache, fullOp), debugString); + return createOperation(cache, createOp(cache, fullOp), debugString, + locInfo); } if (auto absOp = dyn_cast(op); absOp) { - return createOperation(cache, createEltwiseOp(cache, absOp), debugString); + return createOperation(cache, createEltwiseOp(cache, absOp), debugString, + locInfo); } if (auto addOp = dyn_cast(op); addOp) { - return createOperation(cache, createEltwiseOp(cache, addOp), debugString); + return createOperation(cache, createEltwiseOp(cache, addOp), debugString, + locInfo); } if (auto floorOp = dyn_cast(op); floorOp) { - return createOperation(cache, createEltwiseOp(cache, floorOp), debugString); + return createOperation(cache, createEltwiseOp(cache, floorOp), debugString, + locInfo); } if (auto isFiniteOp = dyn_cast(op); isFiniteOp) { return createOperation(cache, createEltwiseOp(cache, isFiniteOp), - debugString); + debugString, locInfo); } if (auto andOp = dyn_cast(op); andOp) { - return createOperation(cache, createEltwiseOp(cache, andOp), debugString); + return createOperation(cache, createEltwiseOp(cache, andOp), debugString, + locInfo); } if (auto cbrtOp = dyn_cast(op); cbrtOp) { - return createOperation(cache, createEltwiseOp(cache, cbrtOp), debugString); + return createOperation(cache, createEltwiseOp(cache, cbrtOp), debugString, + locInfo); } if (auto notOp = dyn_cast(op); notOp) { - return createOperation(cache, createEltwiseOp(cache, notOp), debugString); + return createOperation(cache, createEltwiseOp(cache, notOp), debugString, + locInfo); } if (auto orOp = dyn_cast(op); orOp) { - return createOperation(cache, createEltwiseOp(cache, orOp), debugString); + return createOperation(cache, createEltwiseOp(cache, orOp), debugString, + locInfo); } if (auto xorOp = dyn_cast(op); xorOp) { - return createOperation(cache, createEltwiseOp(cache, xorOp), debugString); + return createOperation(cache, createEltwiseOp(cache, xorOp), debugString, + locInfo); } if (auto multiplyOp = dyn_cast(op); multiplyOp) { return createOperation(cache, createEltwiseOp(cache, multiplyOp), - debugString); + debugString, locInfo); } if (auto negOp = dyn_cast(op); negOp) { - return createOperation(cache, createEltwiseOp(cache, negOp), debugString); + return createOperation(cache, createEltwiseOp(cache, negOp), debugString, + locInfo); } if (auto subtractOp = dyn_cast(op); subtractOp) { return createOperation(cache, createEltwiseOp(cache, subtractOp), - debugString); + debugString, locInfo); } if (auto eqOp = dyn_cast(op); eqOp) { - return createOperation(cache, createEltwiseOp(cache, eqOp), debugString); + return createOperation(cache, createEltwiseOp(cache, eqOp), debugString, + locInfo); } if (auto neOp = dyn_cast(op); neOp) { - return createOperation(cache, createEltwiseOp(cache, neOp), debugString); + return createOperation(cache, createEltwiseOp(cache, neOp), debugString, + locInfo); } if (auto geOp = dyn_cast(op); geOp) { - return createOperation(cache, createEltwiseOp(cache, geOp), debugString); + return createOperation(cache, createEltwiseOp(cache, geOp), debugString, + locInfo); } if (auto gtOp = dyn_cast(op); gtOp) { - return createOperation(cache, createEltwiseOp(cache, gtOp), debugString); + return createOperation(cache, createEltwiseOp(cache, gtOp), debugString, + locInfo); } if (auto leOp = dyn_cast(op); leOp) { - return createOperation(cache, createEltwiseOp(cache, leOp), debugString); + return createOperation(cache, createEltwiseOp(cache, leOp), debugString, + locInfo); } if (auto ltOp = dyn_cast(op); ltOp) { - return createOperation(cache, createEltwiseOp(cache, ltOp), debugString); + return createOperation(cache, createEltwiseOp(cache, ltOp), debugString, + locInfo); } if (auto maximumOp = dyn_cast(op); maximumOp) { return createOperation(cache, createEltwiseOp(cache, maximumOp), - debugString); + debugString, locInfo); } if (auto minimumOp = dyn_cast(op); minimumOp) { return createOperation(cache, createEltwiseOp(cache, minimumOp), - debugString); + debugString, locInfo); } if (auto reluOp = dyn_cast(op); reluOp) { - return createOperation(cache, createEltwiseOp(cache, reluOp), debugString); + return createOperation(cache, createEltwiseOp(cache, reluOp), debugString, + locInfo); } if (auto sqrtOp = dyn_cast(op); sqrtOp) { - return createOperation(cache, createEltwiseOp(cache, sqrtOp), debugString); + return createOperation(cache, createEltwiseOp(cache, sqrtOp), debugString, + locInfo); } if (auto rsqrtOp = dyn_cast(op); rsqrtOp) { - return createOperation(cache, createEltwiseOp(cache, rsqrtOp), debugString); + return createOperation(cache, createEltwiseOp(cache, rsqrtOp), debugString, + locInfo); } if (auto signOp = dyn_cast(op); signOp) { - return createOperation(cache, createEltwiseOp(cache, signOp), debugString); + return createOperation(cache, createEltwiseOp(cache, signOp), debugString, + locInfo); } if (auto expOp = dyn_cast(op); expOp) { - return createOperation(cache, createEltwiseOp(cache, expOp), debugString); + return createOperation(cache, createEltwiseOp(cache, expOp), debugString, + locInfo); } if (auto logOp = dyn_cast(op); logOp) { - return createOperation(cache, createEltwiseOp(cache, logOp), debugString); + return createOperation(cache, createEltwiseOp(cache, logOp), debugString, + locInfo); } if (auto expm1Op = dyn_cast(op); expm1Op) { - return createOperation(cache, createEltwiseOp(cache, expm1Op), debugString); + return createOperation(cache, createEltwiseOp(cache, expm1Op), debugString, + locInfo); } if (auto sigmoidOp = dyn_cast(op); sigmoidOp) { return createOperation(cache, createEltwiseOp(cache, sigmoidOp), - debugString); + debugString, locInfo); } if (auto log1pOp = dyn_cast(op); log1pOp) { - return createOperation(cache, createEltwiseOp(cache, log1pOp), debugString); + return createOperation(cache, createEltwiseOp(cache, log1pOp), debugString, + locInfo); } if (auto scatterOp = dyn_cast(op); scatterOp) { return createOperation(cache, createEltwiseOp(cache, scatterOp), - debugString); + debugString, locInfo); } if (auto reciprocalOp = dyn_cast(op); reciprocalOp) { return createOperation(cache, createEltwiseOp(cache, reciprocalOp), - debugString); + debugString, locInfo); } if (auto divOp = dyn_cast(op); divOp) { - return createOperation(cache, createEltwiseOp(cache, divOp), debugString); + return createOperation(cache, createEltwiseOp(cache, divOp), debugString, + locInfo); } if (auto remainderOp = dyn_cast(op); remainderOp) { return createOperation(cache, createEltwiseOp(cache, remainderOp), - debugString); + debugString, locInfo); } if (auto leakyReluOp = dyn_cast(op); leakyReluOp) { return createOperation(cache, createEltwiseOp(cache, leakyReluOp), - debugString); + debugString, locInfo); } if (auto linearOp = dyn_cast(op); linearOp) { - return createOperation(cache, createOp(cache, linearOp), debugString); + return createOperation(cache, createOp(cache, linearOp), debugString, + locInfo); } if (auto matmulOp = dyn_cast(op); matmulOp) { - return createOperation(cache, createOp(cache, matmulOp), debugString); + return createOperation(cache, createOp(cache, matmulOp), debugString, + locInfo); } if (auto sumOp = dyn_cast(op); sumOp) { - return createOperation(cache, createReductionOp(cache, sumOp), debugString); + return createOperation(cache, createReductionOp(cache, sumOp), debugString, + locInfo); } if (auto meanOp = dyn_cast(op); meanOp) { - return createOperation(cache, createReductionOp(cache, meanOp), - debugString); + return createOperation(cache, createReductionOp(cache, meanOp), debugString, + locInfo); } if (auto maxOp = dyn_cast(op); maxOp) { - return createOperation(cache, createReductionOp(cache, maxOp), debugString); + return createOperation(cache, createReductionOp(cache, maxOp), debugString, + locInfo); } if (auto embeddingOp = dyn_cast(op); embeddingOp) { return createOperation(cache, createEmbeddingOp(cache, embeddingOp), - debugString); + debugString, locInfo); } if (auto softmaxOp = dyn_cast(op); softmaxOp) { return createOperation(cache, createSoftmaxOp(cache, softmaxOp), - debugString); + debugString, locInfo); } if (auto transposeOp = dyn_cast(op); transposeOp) { return createOperation(cache, createTransposeOp(cache, transposeOp), - debugString); + debugString, locInfo); } if (auto clampOp = dyn_cast(op); clampOp) { return createOperation(cache, createNonDPSEltwiseOp(cache, clampOp), - debugString); + debugString, locInfo); } if (auto conv2dOp = dyn_cast(op); conv2dOp) { - return createOperation(cache, createOp(cache, conv2dOp), debugString); + return createOperation(cache, createOp(cache, conv2dOp), debugString, + locInfo); } if (auto allGatherOp = dyn_cast(op); allGatherOp) { - return createOperation(cache, createOp(cache, allGatherOp), debugString); + return createOperation(cache, createOp(cache, allGatherOp), debugString, + locInfo); } if (auto concatOp = dyn_cast(op); concatOp) { - return createOperation(cache, createConcatOp(cache, concatOp), debugString); + return createOperation(cache, createConcatOp(cache, concatOp), debugString, + locInfo); } if (auto reshapeOp = dyn_cast(op); reshapeOp) { return createOperation(cache, createReshapeOp(cache, reshapeOp), - debugString); + debugString, locInfo); } if (auto sliceOp = dyn_cast(op); sliceOp) { - return createOperation(cache, createSliceOp(cache, sliceOp), debugString); + return createOperation(cache, createSliceOp(cache, sliceOp), debugString, + locInfo); } if (auto max_pool2dOp = dyn_cast(op); max_pool2dOp) { return createOperation(cache, createMaxPool2dOp(cache, max_pool2dOp), - debugString); + debugString, locInfo); } if (auto deallocateOp = dyn_cast(op); deallocateOp) { return createOperation(cache, createDeallocateOp(cache, deallocateOp), - debugString); + debugString, locInfo); } if (auto ceilOp = dyn_cast(op); ceilOp) { - return createOperation(cache, createEltwiseOp(cache, ceilOp), debugString); + return createOperation(cache, createEltwiseOp(cache, ceilOp), debugString, + locInfo); } if (auto cosOp = dyn_cast(op); cosOp) { - return createOperation(cache, createEltwiseOp(cache, cosOp), debugString); + return createOperation(cache, createEltwiseOp(cache, cosOp), debugString, + locInfo); } if (auto sinOp = dyn_cast(op); sinOp) { - return createOperation(cache, createEltwiseOp(cache, sinOp), debugString); + return createOperation(cache, createEltwiseOp(cache, sinOp), debugString, + locInfo); } if (auto whereOp = dyn_cast(op); whereOp) { - return createOperation(cache, createEltwiseOp(cache, whereOp), debugString); + return createOperation(cache, createEltwiseOp(cache, whereOp), debugString, + locInfo); } if (auto geluOp = dyn_cast(op); geluOp) { - return createOperation(cache, createEltwiseOp(cache, geluOp), debugString); + return createOperation(cache, createEltwiseOp(cache, geluOp), debugString, + locInfo); } if (auto arangeOp = dyn_cast(op); arangeOp) { - return createOperation(cache, createOp(cache, arangeOp), debugString); + return createOperation(cache, createOp(cache, arangeOp), debugString, + locInfo); } llvm_unreachable("unhandled op in emitTTNNOperation"); diff --git a/runtime/include/tt/runtime/detail/ttmetal.h b/runtime/include/tt/runtime/detail/ttmetal.h index 7a68a7e94..5544e1d70 100644 --- a/runtime/include/tt/runtime/detail/ttmetal.h +++ b/runtime/include/tt/runtime/detail/ttmetal.h @@ -47,6 +47,8 @@ void wait(Event event); std::string getOpDebugString(OpContext opContextHandle); +std::string getOpLocInfo(OpContext opContextHandle); + Tensor getOpOutputTensor(OpContext opContextHandle, CallbackContext programContextHandle); diff --git a/runtime/include/tt/runtime/detail/ttnn.h b/runtime/include/tt/runtime/detail/ttnn.h index 6c55ac1de..67aa91a71 100644 --- a/runtime/include/tt/runtime/detail/ttnn.h +++ b/runtime/include/tt/runtime/detail/ttnn.h @@ -83,6 +83,8 @@ void wait(Event event); std::string getOpDebugString(OpContext opContextHandle); +std::string getOpLocInfo(OpContext opContextHandle); + Tensor getOpOutputTensor(OpContext opContextHandle, CallbackContext programContextHandle); diff --git a/runtime/include/tt/runtime/runtime.h b/runtime/include/tt/runtime/runtime.h index 1dc721f66..e4348da60 100644 --- a/runtime/include/tt/runtime/runtime.h +++ b/runtime/include/tt/runtime/runtime.h @@ -71,6 +71,8 @@ void wait(Event event); std::string getOpDebugString(OpContext opContextHandle); +std::string getOpLocInfo(OpContext opContextHandle); + Tensor getOpOutputTensor(OpContext opContextHandle, CallbackContext programContextHandle); diff --git a/runtime/lib/runtime.cpp b/runtime/lib/runtime.cpp index 586b8394e..a57ac3fcd 100644 --- a/runtime/lib/runtime.cpp +++ b/runtime/lib/runtime.cpp @@ -261,6 +261,21 @@ std::string getOpDebugString(OpContext opContextHandle) { throw std::runtime_error("runtime is not enabled"); } +std::string getOpLocInfo(OpContext opContextHandle) { +#ifdef TT_RUNTIME_ENABLE_TTNN + if (getCurrentRuntime() == DeviceRuntime::TTNN) { + return ::tt::runtime::ttnn::getOpLocInfo(opContextHandle); + } +#endif + +#ifdef TT_RUNTIME_ENABLE_TTMETAL + if (getCurrentRuntime() == DeviceRuntime::TTMetal) { + return ::tt::runtime::ttmetal::getOpLocInfo(opContextHandle); + } +#endif + throw std::runtime_error("runtime is not enabled"); +} + Tensor getOpOutputTensor(OpContext opContextHandle, CallbackContext programContextHandle) { #if defined(TT_RUNTIME_ENABLE_TTNN) diff --git a/runtime/lib/ttmetal/runtime.cpp b/runtime/lib/ttmetal/runtime.cpp index ab343554e..22d43ba36 100644 --- a/runtime/lib/ttmetal/runtime.cpp +++ b/runtime/lib/ttmetal/runtime.cpp @@ -262,6 +262,12 @@ std::string getOpDebugString(OpContext opContextHandle) { return ""; } +std::string getOpLocInfo(OpContext opContextHandle) { + // Not implemented + LOG_WARNING("obtaining op location info for metal runtime not implemented"); + return ""; +} + Tensor getOpOutputTensor(OpContext opContextHandle, CallbackContext programContextHandle) { // Not implemented diff --git a/runtime/lib/ttnn/program.cpp b/runtime/lib/ttnn/program.cpp index 48b0be7ff..3aab3a94c 100644 --- a/runtime/lib/ttnn/program.cpp +++ b/runtime/lib/ttnn/program.cpp @@ -33,9 +33,19 @@ #include "tt/runtime/utils.h" #include "ttmlir/Target/TTNN/program_generated.h" +#ifdef TT_RUNTIME_ENABLE_PERF_TRACE +#include "tracy/Tracy.hpp" +#endif + namespace tt::runtime::ttnn { using LogType = ::tt::runtime::logger::LogType; +void tracyLogOpLocation(const ::tt::target::ttnn::Operation *op) { +#ifdef TT_RUNTIME_ENABLE_PERF_TRACE + TracyMessage(op->loc_info()->c_str(), op->loc_info()->size()); +#endif +} + static ::tt::target::ttnn::TTNNBinary const *getBinary(Flatbuffer binary) { bool isTTNN = ::tt::target::ttnn::SizePrefixedTTNNBinaryBufferHasIdentifier( binary.handle.get()); @@ -74,6 +84,7 @@ class ProgramExecutor { for (const ::tt::target::ttnn::Operation *op : *program->operations()) { LOG_DEBUG(LogType::LogRuntimeTTNN, "Executing operation: ", op->debug_info()->c_str()); + tracyLogOpLocation(op); runOperation(op); runCallback(executableHandle, op, &context); } diff --git a/runtime/lib/ttnn/runtime.cpp b/runtime/lib/ttnn/runtime.cpp index 86fd2d25c..2dfc07788 100644 --- a/runtime/lib/ttnn/runtime.cpp +++ b/runtime/lib/ttnn/runtime.cpp @@ -202,6 +202,12 @@ std::string getOpDebugString(OpContext opContextHandle) { return std::string(opContext.debug_info()->c_str()); } +std::string getOpLocInfo(OpContext opContextHandle) { + auto const &opContext = + opContextHandle.as<::tt::target::ttnn::Operation>(DeviceRuntime::TTNN); + return std::string(opContext.loc_info()->c_str()); +} + Tensor getOpOutputTensor(OpContext opContextHandle, CallbackContext programContextHandle) { auto const &programContext = diff --git a/runtime/tools/python/ttrt/common/perf.py b/runtime/tools/python/ttrt/common/perf.py index a341c2b4f..f70defa31 100644 --- a/runtime/tools/python/ttrt/common/perf.py +++ b/runtime/tools/python/ttrt/common/perf.py @@ -17,11 +17,16 @@ import atexit import traceback from pathlib import Path +import csv from ttrt.common.util import * from ttrt.common.query import Query +def get_loc_data_hook(binary, programContext, opContext): + op_debug_str = ttrt.runtime.get_op_debug_str(opContext) + + class Perf: registered_args = {} @@ -456,6 +461,38 @@ def signal_handler(sig, frame): ) process_ops(None, None, False) + + # Add post-processing steps to insert location data into the ops_perf data file + with open(profiler_csv_file_path, "r") as perf_file: + perf_reader = csv.DictReader(perf_file) + headers = list(perf_reader.fieldnames) + ["LOC"] + perf_data = list(perf_reader) + + with open(profiler_csv_file_path, "w+") as perf_file, open( + tracy_ops_data_file_path, "r" + ) as message_file: + message_reader = csv.reader(message_file, delimiter=";") + ops_index = 0 + prev = None + for message in message_reader: + message = message[0] # Don't need timestamp information + if message.startswith("`"): + # This is a TTNN Message + # The location data is now in the previous message + # The order of data is maintained in perf_data so as the messages are received, they update the id last encountered. + # Now that we have a new message, we can update the location data from the previous message + if prev: + # Get the location data from the previous message and add it as new data for the perf_data (as a new col) + if len(perf_data) > ops_index: + perf_data[ops_index]["LOC"] = prev + ops_index += 1 + else: + prev = message + perf_writer = csv.DictWriter(perf_file, fieldnames=headers) + perf_writer.writeheader() + for row in perf_data: + perf_writer.writerow(row) + self.file_manager.copy_file( perf_folder_path, profiler_csv_file_path, diff --git a/runtime/tools/python/ttrt/runtime/module.cpp b/runtime/tools/python/ttrt/runtime/module.cpp index dfc4a6820..c0378727c 100644 --- a/runtime/tools/python/ttrt/runtime/module.cpp +++ b/runtime/tools/python/ttrt/runtime/module.cpp @@ -100,6 +100,8 @@ PYBIND11_MODULE(_C, m) { "Get the input tensor of the op"); m.def("get_op_debug_str", &tt::runtime::getOpDebugString, "Get the debug string of the op"); + m.def("get_op_loc_info", &tt::runtime::getOpLocInfo, + "Get the location info of the op"); py::class_(m, "DebugEnv") .def_static("get", &tt::runtime::debug::Env::get) diff --git a/third_party/CMakeLists.txt b/third_party/CMakeLists.txt index e033913e2..90173880d 100644 --- a/third_party/CMakeLists.txt +++ b/third_party/CMakeLists.txt @@ -28,6 +28,7 @@ set(TTMETAL_INCLUDE_DIRS ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/hw/inc/${ARCH_EXTRA_DIR} ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_metal/third_party/umd/src/firmware/riscv/${ARCH_NAME} ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/tt_eager + ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal-build/include ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/reflect/e75434c4c5f669e4a74e4d84e0a30d7249c1e66f ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/nanomsg/28cc32d5bdb6a858fe53b3ccf7e923957e53eada/include ${PROJECT_SOURCE_DIR}/third_party/tt-metal/src/tt-metal/.cpmcache/fmt/73b5ec45edbd92babfd91c3777a9e1ab9cac8238/include