Star119
54
代码
代码
Issues24
Pull Requests26
项目讨论
Wiki
分析
Star119
54
已合并
支持解析launchKernel接口上报额外的profiling信息 #252
xfeng创建于 5月26日
支持解析launchKernel接口上报额外的profiling信息 #252
已合并
xfeng创建于 5月26日
从zyb_230/msprofmain合入到Ascend/msprofmaster
共 22 个文件变更+1746-1061
Manalysis/csrc/application/database/db_assembler.cpp+6-2
@@ -71,7 +71,7 @@ const std::vector<std::string> COMM_TASK_INDEX_COLS = {"globalTaskId"};
71using CommScheduleDataFormat = std::vector<std::tuple<uint64_t, uint64_t, uint64_t, uint64_t>>;71using CommScheduleDataFormat = std::vector<std::tuple<uint64_t, uint64_t, uint64_t, uint64_t>>;
72using ComputeTaskInfoFormat =72using ComputeTaskInfoFormat =
73 std::vector<std::tuple<uint64_t, uint64_t, uint32_t, uint32_t, uint64_t, uint64_t, uint64_t, uint64_t, uint64_t,73 std::vector<std::tuple<uint64_t, uint64_t, uint32_t, uint32_t, uint64_t, uint64_t, uint64_t, uint64_t, uint64_t,
74- uint64_t, uint64_t, uint64_t, uint64_t, uint64_t, uint64_t>>;74+ uint64_t, uint64_t, uint64_t, uint64_t, uint64_t, uint64_t, uint64_t, uint64_t>>;
75// 大算子数据75// 大算子数据
76// opName, start, end, connectionId, group_name, opId, relay, retry, data_type, alg_type, count, op_type, deviceId,76// opName, start, end, connectionId, group_name, opId, relay, retry, data_type, alg_type, count, op_type, deviceId,
77// rank_size77// rank_size
@@ -100,6 +100,8 @@ struct ComputeTaskInfoData
100 uint64_t hashId; // 对应attrInfo100 uint64_t hashId; // 对应attrInfo
101 uint64_t opState;101 uint64_t opState;
102 uint64_t hf32Eligible; // opFlag102 uint64_t hf32Eligible; // opFlag
103+ uint64_t gridDim;
104+ uint64_t blockDim;
103};105};
104 106 
105std::string ReplaceQuotes(const std::string& input)107std::string ReplaceQuotes(const std::string& input)
@@ -826,11 +828,13 @@ bool SaveComputeTaskInfo(DataInventory& dataInventory, DBInfo& msprofDB, const s
826 taskInfoData.hashId = IdPool::GetInstance().GetUint64Id(item.hashId);828 taskInfoData.hashId = IdPool::GetInstance().GetUint64Id(item.hashId);
827 taskInfoData.opState = IdPool::GetInstance().GetUint64Id(item.opState);829 taskInfoData.opState = IdPool::GetInstance().GetUint64Id(item.opState);
828 taskInfoData.hf32Eligible = IdPool::GetInstance().GetUint64Id(item.opFlag);830 taskInfoData.hf32Eligible = IdPool::GetInstance().GetUint64Id(item.opFlag);
831+ taskInfoData.gridDim = IdPool::GetInstance().GetUint64Id(item.gridDim);
832+ taskInfoData.blockDim = IdPool::GetInstance().GetUint64Id(item.blockDim);
829 res.emplace_back(taskInfoData.opName, taskInfoData.globalTaskId, item.blockNum, item.mixBlockNum,833 res.emplace_back(taskInfoData.opName, taskInfoData.globalTaskId, item.blockNum, item.mixBlockNum,
830 taskInfoData.taskType, taskInfoData.opType, taskInfoData.inputFormats,834 taskInfoData.taskType, taskInfoData.opType, taskInfoData.inputFormats,
831 taskInfoData.inputDataTypes, taskInfoData.inputShapes, taskInfoData.outputFormats,835 taskInfoData.inputDataTypes, taskInfoData.inputShapes, taskInfoData.outputFormats,
832 taskInfoData.outputDataTypes, taskInfoData.outputShapes, taskInfoData.hashId,836 taskInfoData.outputDataTypes, taskInfoData.outputShapes, taskInfoData.hashId,
833- taskInfoData.opState, taskInfoData.hf32Eligible);837+ taskInfoData.opState, taskInfoData.hf32Eligible, taskInfoData.gridDim, taskInfoData.blockDim);
834 }838 }
835 bool flag = true;839 bool flag = true;
836 if (!res.empty())840 if (!res.empty())
Manalysis/csrc/application/timeline/ascend_hardware_assembler.cpp+146-85
@@ -15,51 +15,51 @@
15 * -------------------------------------------------------------------------*/15 * -------------------------------------------------------------------------*/
16 16 
17#include "analysis/csrc/application/timeline/ascend_hardware_assembler.h"17#include "analysis/csrc/application/timeline/ascend_hardware_assembler.h"
18+ 
18#include "analysis/csrc/application//credential/id_pool.h"19#include "analysis/csrc/application//credential/id_pool.h"
19#include "analysis/csrc/application/timeline/connection_id_pool.h"20#include "analysis/csrc/application/timeline/connection_id_pool.h"
20#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/api_data.h"21#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/api_data.h"
21#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/memcpy_info_data.h"22#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/memcpy_info_data.h"
22 23 
23-namespace Analysis {24+namespace Analysis
24-namespace Application {25+{
26+namespace Application
27+{
25using namespace Analysis::Viewer::Database;28using namespace Analysis::Viewer::Database;
26using namespace Analysis::Utils;29using namespace Analysis::Utils;
27using namespace Analysis::Domain;30using namespace Analysis::Domain;
28using IdPool = Analysis::Application::Credential::IdPool;31using IdPool = Analysis::Application::Credential::IdPool;
29-namespace {32+namespace
33+{
30using MEMCPY_INFO_FORMAT = std::map<TaskId, MemcpyInfoData>;34using MEMCPY_INFO_FORMAT = std::map<TaskId, MemcpyInfoData>;
31const std::string TASK_TYPE_FFTS_PLUS = "FFTS_PLUS";35const std::string TASK_TYPE_FFTS_PLUS = "FFTS_PLUS";
32const std::string TASK_TYPE_UNKNOWN = "UNKNOWN";36const std::string TASK_TYPE_UNKNOWN = "UNKNOWN";
33const std::string TASK_TYPE_NA = "N/A";37const std::string TASK_TYPE_NA = "N/A";
34-const std::vector<std::string> MEMCPY_OPERATIONS {38+const std::vector<std::string> MEMCPY_OPERATIONS{"host to host", "host to device", "device to host",
35- "host to host",39+ "device to device", "managed memory", "addr device to device",
36- "host to device",40+ "host to device ex", "device to host ex"};
37- "device to host",
38- "device to device",
39- "managed memory",
40- "addr device to device",
41- "host to device ex",
42- "device to host ex"
43-};
44const std::string OTHER_DIRECTION = "other";41const std::string OTHER_DIRECTION = "other";
45-}42+} // namespace
46 43 
47MEMCPY_INFO_FORMAT GenerateMemcpyInfoDataMap(const std::shared_ptr<std::vector<MemcpyInfoData>> &res)44MEMCPY_INFO_FORMAT GenerateMemcpyInfoDataMap(const std::shared_ptr<std::vector<MemcpyInfoData>> &res)
48{45{
49 MEMCPY_INFO_FORMAT memcpyInfoDataMap;46 MEMCPY_INFO_FORMAT memcpyInfoDataMap;
50- if (res != nullptr) {47+ if (res != nullptr)
51- for (const auto &item: *res) {48+ {
49+ for (const auto &item : *res)
50+ {
52 memcpyInfoDataMap[item.taskId] = std::move(item);51 memcpyInfoDataMap[item.taskId] = std::move(item);
53 }52 }
54 }53 }
55 return memcpyInfoDataMap;54 return memcpyInfoDataMap;
56}55}
57 56 
58- 
59AscendHardwareAssembler::AscendHardwareAssembler()57AscendHardwareAssembler::AscendHardwareAssembler()
60- : JsonAssembler(PROCESS_TASK, {{MSPROF_JSON_FILE, FileCategory::MSPROF}}) {}58+ : JsonAssembler(PROCESS_TASK, {{MSPROF_JSON_FILE, FileCategory::MSPROF}})
59+{
60+}
61 61 
62-void TaskTraceEvent::ProcessArgs(JsonWriter& ostream)62+void TaskTraceEvent::ProcessArgs(JsonWriter &ostream)
63{63{
64 ostream["Model Id"] << modelId_;64 ostream["Model Id"] << modelId_;
65 ostream["Task Type"] << taskType_;65 ostream["Task Type"] << taskType_;
@@ -70,17 +70,25 @@ void TaskTraceEvent::ProcessArgs(JsonWriter& ostream)
70 ostream["connection_id"] << connectionId_;70 ostream["connection_id"] << connectionId_;
71}71}
72 72 
73-void MemcpyAsyncEvent::ProcessArgs(JsonWriter& ostream)73+void MemcpyAsyncEvent::ProcessArgs(JsonWriter &ostream)
74{74{
75 TaskTraceEvent::ProcessArgs(ostream);75 TaskTraceEvent::ProcessArgs(ostream);
76- if (showFlag_) {76+ if (showFlag_)
77+ {
77 ostream["size(B)"] << dataSize_;78 ostream["size(B)"] << dataSize_;
78 ostream["bandwidth(GB/s)"] << bandwidth_;79 ostream["bandwidth(GB/s)"] << bandwidth_;
79 ostream["operation"] << memcpyDirection_;80 ostream["operation"] << memcpyDirection_;
80 }81 }
81}82}
82 83 
83-void KfcTurnTraceEvent::ProcessArgs(JsonWriter& ostream)84+void SimtTaskEvent::ProcessArgs(JsonWriter &ostream)
85+{
86+ TaskTraceEvent::ProcessArgs(ostream);
87+ ostream["Grid Dim"] << gridDim_;
88+ ostream["Block Dim"] << blockDim_;
89+}
90+ 
91+void KfcTurnTraceEvent::ProcessArgs(JsonWriter &ostream)
84{92{
85 ostream["Physic Stream Id"] << streamId_;93 ostream["Physic Stream Id"] << streamId_;
86 ostream["Task Id"] << taskId_;94 ostream["Task Id"] << taskId_;
@@ -90,53 +98,65 @@ void AscendHardwareAssembler::InitData(DataInventory &dataInventory, std::vector
90{98{
91 logicStream_ = dataInventory.GetPtr<std::unordered_map<uint32_t, uint32_t>>();99 logicStream_ = dataInventory.GetPtr<std::unordered_map<uint32_t, uint32_t>>();
92 auto taskInfo = dataInventory.GetPtr<std::vector<TaskInfoData>>();100 auto taskInfo = dataInventory.GetPtr<std::vector<TaskInfoData>>();
93- if (taskInfo != nullptr) {101+ if (taskInfo != nullptr)
94- for (const auto &node : *taskInfo) {102+ {
95- const TaskId& taskId = TaskId{static_cast<uint16_t >(node.streamId), static_cast<uint16_t >(node.batchId),103+ for (const auto &node : *taskInfo)
96- node.taskId, node.contextId, node.deviceId};104+ {
105+ const TaskId &taskId = TaskId{static_cast<uint16_t>(node.streamId), static_cast<uint16_t>(node.batchId),
106+ node.taskId, node.contextId, node.deviceId};
97 opName_.emplace(taskId, node.opName);107 opName_.emplace(taskId, node.opName);
98 taskType_.emplace(taskId, node.taskType);108 taskType_.emplace(taskId, node.taskType);
109+ simtInfoMap_.emplace(taskId, std::make_pair(node.gridDim, node.blockDim));
99 }110 }
100 }111 }
101 auto apiData = dataInventory.GetPtr<std::vector<ApiData>>();112 auto apiData = dataInventory.GetPtr<std::vector<ApiData>>();
102- if (apiData != nullptr) {113+ if (apiData != nullptr)
103- for (const auto &node : *apiData) {114+ {
104- if (RECORD_EVENT == node.id || WAIT_EVENT == node.id) {115+ for (const auto &node : *apiData)
116+ {
117+ if (RECORD_EVENT == node.id || WAIT_EVENT == node.id)
118+ {
105 aclEvent_.emplace(node.connectionId);119 aclEvent_.emplace(node.connectionId);
106 }120 }
107 }121 }
108 }122 }
109- for (const auto& data : taskData) {123+ for (const auto &data : taskData)
110- if (data.contextId != UINT32_MAX) {124+ {
111- ffts_.emplace(TaskId{static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId),125+ if (data.contextId != UINT32_MAX)
112- data.taskId, UINT32_MAX, data.deviceId});126+ {
127+ ffts_.emplace(TaskId{static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId), data.taskId,
128+ UINT32_MAX, data.deviceId});
113 }129 }
114- if (data.hostType == MEMCPY_ASYNC) {130+ if (data.hostType == MEMCPY_ASYNC)
131+ {
115 memcpyAsyncDeviceTasks_.push_back(data);132 memcpyAsyncDeviceTasks_.push_back(data);
116 }133 }
117 }134 }
118}135}
119 136 
120-std::string AscendHardwareAssembler::GetOpName(const AscendTaskData& data)137+std::string AscendHardwareAssembler::GetOpName(const AscendTaskData &data)
121{138{
122- TaskId id{static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId),139+ TaskId id{static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId), data.taskId, data.contextId,
123- data.taskId, data.contextId, data.deviceId};140+ data.deviceId};
124 auto it = opName_.find(id);141 auto it = opName_.find(id);
125- if (it != opName_.end()) {142+ if (it != opName_.end())
143+ {
126 return it->second;144 return it->second;
127 }145 }
128- if (data.hostType == TASK_TYPE_FFTS_PLUS || data.hostType == TASK_TYPE_UNKNOWN) {146+ if (data.hostType == TASK_TYPE_FFTS_PLUS || data.hostType == TASK_TYPE_UNKNOWN)
147+ {
129 return data.deviceType;148 return data.deviceType;
130 }149 }
131 return data.hostType;150 return data.hostType;
132}151}
133 152 
134-std::string AscendHardwareAssembler::GetTaskType(const AscendTaskData& data)153+std::string AscendHardwareAssembler::GetTaskType(const AscendTaskData &data)
135{154{
136- TaskId id{static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId),155+ TaskId id{static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId), data.taskId, data.contextId,
137- data.taskId, data.contextId, data.deviceId};156+ data.deviceId};
138 auto it = taskType_.find(id);157 auto it = taskType_.find(id);
139- if (it != taskType_.end() && it->second != TASK_TYPE_NA) {158+ if (it != taskType_.end() && it->second != TASK_TYPE_NA)
159+ {
140 return it->second;160 return it->second;
141 }161 }
142 return data.taskType;162 return data.taskType;
@@ -144,11 +164,13 @@ std::string AscendHardwareAssembler::GetTaskType(const AscendTaskData& data)
144 164 
145uint32_t AscendHardwareAssembler::GetPhysicStreamId(const uint32_t streamId)165uint32_t AscendHardwareAssembler::GetPhysicStreamId(const uint32_t streamId)
146{166{
147- if (logicStream_ == nullptr) {167+ if (logicStream_ == nullptr)
168+ {
148 return streamId;169 return streamId;
149 }170 }
150 auto it = logicStream_->find(streamId);171 auto it = logicStream_->find(streamId);
151- if (it != logicStream_->end()) {172+ if (it != logicStream_->end())
173+ {
152 return it->second;174 return it->second;
153 }175 }
154 return streamId;176 return streamId;
@@ -162,27 +184,49 @@ void AscendHardwareAssembler::GenerateTaskTrace(const std::vector<AscendTaskData
162 std::string traceName;184 std::string traceName;
163 std::string taskTypeName;185 std::string taskTypeName;
164 TaskId id;186 TaskId id;
165- for (const auto &data : taskData) {187+ for (const auto &data : taskData)
166- if (data.hostType == MEMCPY_ASYNC) {188+ {
189+ if (data.hostType == MEMCPY_ASYNC)
190+ {
167 continue; // MEMCPY_ASYNC类型的task有新增args,需要单独处理191 continue; // MEMCPY_ASYNC类型的task有新增args,需要单独处理
168 }192 }
169- id = {static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId),193+ id = {static_cast<uint16_t>(data.streamId), static_cast<uint16_t>(data.batchId), data.taskId, data.contextId,
170- data.taskId, data.contextId, data.deviceId};194+ data.deviceId};
171- if (ffts_.find(id) != ffts_.end()) { // 当前task存在ffts+任务，只呈现ffts+任务即可195+ if (ffts_.find(id) != ffts_.end())
196+ { // 当前task存在ffts+任务，只呈现ffts+任务即可
172 continue;197 continue;
173 }198 }
174 traceName = GetOpName(data);199 traceName = GetOpName(data);
175 taskTypeName = GetTaskType(data);200 taskTypeName = GetTaskType(data);
176 formatPid = GetDevicePid(pidMap, data.deviceId, profPath, layer.sortIndex);201 formatPid = GetDevicePid(pidMap, data.deviceId, profPath, layer.sortIndex);
177 int tid = static_cast<int>(GetPhysicStreamId(data.streamId));202 int tid = static_cast<int>(GetPhysicStreamId(data.streamId));
178- // 存储pid，tid组合的最小集
179 pidTidSet_.insert({formatPid, tid});203 pidTidSet_.insert({formatPid, tid});
180- std::shared_ptr<TaskTraceEvent> event;204+ if (data.taskType == KERNEL_SIMT_TASK_TYPE)
181- MAKE_SHARED_RETURN_VOID(event, TaskTraceEvent, formatPid, tid, data.duration / NS_TO_US,205+ {
182- DivideByPowersOfTenWithPrecision(data.timestamp), traceName,206+ std::string gridDim = NA;
183- data.modelId, data.streamId,207+ std::string blockDim = NA;
184- data.taskId, data.batchId, data.contextId, data.connectionId, taskTypeName);208+ auto it = simtInfoMap_.find(id);
185- res_.push_back(event);209+ if (it != simtInfoMap_.end())
210+ {
211+ gridDim = it->second.first;
Wangang Yu
Wangang YuWangang Yu5月27日

[review] 问题： 直接使用it->second.first和it->second.second，未校验 pair 中的字符串是否为空。 影响：若simtInfoMap_中存在空字符串的条目，会导致gridDim/blockDim为空，影响后续数据展示。 建议：添加空值校验，空值时使用默认值 NA

likedislike
xfeng
xfeng
6月1日 评论：
212+ blockDim = it->second.second;
213+ }
214+ std::shared_ptr<SimtTaskEvent> event;
215+ MAKE_SHARED_RETURN_VOID(event, SimtTaskEvent, formatPid, tid, data.duration / NS_TO_US,
216+ DivideByPowersOfTenWithPrecision(data.timestamp), traceName, data.modelId,
217+ data.streamId, data.taskId, data.batchId, data.contextId, data.connectionId,
218+ taskTypeName, gridDim, blockDim);
219+ res_.push_back(event);
220+ }
221+ else
222+ {
223+ std::shared_ptr<TaskTraceEvent> event;
224+ MAKE_SHARED_RETURN_VOID(event, TaskTraceEvent, formatPid, tid, data.duration / NS_TO_US,
225+ DivideByPowersOfTenWithPrecision(data.timestamp), traceName, data.modelId,
226+ data.streamId, data.taskId, data.batchId, data.contextId, data.connectionId,
227+ taskTypeName);
228+ res_.push_back(event);
229+ }
186 GenerateTaskConnectionTrace(data, formatPid, id);230 GenerateTaskConnectionTrace(data, formatPid, id);
187 }231 }
188}232}
@@ -191,7 +235,8 @@ void AscendHardwareAssembler::GenerateKfcTrace(const std::vector<KfcTurnData> &k
191 const LayerInfo &layer, std::unordered_map<uint16_t, uint32_t> &pidMap)235 const LayerInfo &layer, std::unordered_map<uint16_t, uint32_t> &pidMap)
192{236{
193 uint32_t formatPid;237 uint32_t formatPid;
194- for (const auto &datum: kfcData) {238+ for (const auto &datum : kfcData)
239+ {
195 std::string traceName = datum.opName;240 std::string traceName = datum.opName;
196 formatPid = GetDevicePid(pidMap, datum.deviceId, profPath, layer.sortIndex);241 formatPid = GetDevicePid(pidMap, datum.deviceId, profPath, layer.sortIndex);
197 int formatTid = static_cast<int>(GetPhysicStreamId(datum.streamId));242 int formatTid = static_cast<int>(GetPhysicStreamId(datum.streamId));
@@ -199,14 +244,15 @@ void AscendHardwareAssembler::GenerateKfcTrace(const std::vector<KfcTurnData> &k
199 pidTidSet_.insert({formatPid, formatTid});244 pidTidSet_.insert({formatPid, formatTid});
200 std::shared_ptr<KfcTurnTraceEvent> event;245 std::shared_ptr<KfcTurnTraceEvent> event;
201 MAKE_SHARED_RETURN_VOID(event, KfcTurnTraceEvent, formatPid, formatTid, datum.duration / NS_TO_US,246 MAKE_SHARED_RETURN_VOID(event, KfcTurnTraceEvent, formatPid, formatTid, datum.duration / NS_TO_US,
202- DivideByPowersOfTenWithPrecision(datum.timestamp),247+ DivideByPowersOfTenWithPrecision(datum.timestamp), traceName, datum.streamId,
203- traceName, datum.streamId, datum.taskId);248+ datum.taskId);
204 res_.push_back(event);249 res_.push_back(event);
205 }250 }
206}251}
207 252 
208void AscendHardwareAssembler::GenerateMemcpyAsyncTrace(DataInventory &dataInventory, const std::string &profPath,253void AscendHardwareAssembler::GenerateMemcpyAsyncTrace(DataInventory &dataInventory, const std::string &profPath,
209- const LayerInfo &layer, std::unordered_map<uint16_t, uint32_t> &pidMap)254+ const LayerInfo &layer,
255+ std::unordered_map<uint16_t, uint32_t> &pidMap)
210{256{
211 uint32_t formatPid;257 uint32_t formatPid;
212 std::string traceName;258 std::string traceName;
@@ -215,11 +261,13 @@ void AscendHardwareAssembler::GenerateMemcpyAsyncTrace(DataInventory &dataInvent
215 std::string memcpyDirection;261 std::string memcpyDirection;
216 bool showFlag = true;262 bool showFlag = true;
217 auto memcpyInfo = dataInventory.GetPtr<std::vector<MemcpyInfoData>>();263 auto memcpyInfo = dataInventory.GetPtr<std::vector<MemcpyInfoData>>();
218- if (memcpyInfo == nullptr) {264+ if (memcpyInfo == nullptr)
265+ {
219 showFlag = false;266 showFlag = false;
220 }267 }
221 MEMCPY_INFO_FORMAT memcpyInfoDataMap = GenerateMemcpyInfoDataMap(memcpyInfo);268 MEMCPY_INFO_FORMAT memcpyInfoDataMap = GenerateMemcpyInfoDataMap(memcpyInfo);
222- for (const auto &data : memcpyAsyncDeviceTasks_) {269+ for (const auto &data : memcpyAsyncDeviceTasks_)
270+ {
223 dataSize = 0;271 dataSize = 0;
224 memcpyDirection = OTHER_DIRECTION;272 memcpyDirection = OTHER_DIRECTION;
225 bandwidth = 0.0;273 bandwidth = 0.0;
@@ -229,24 +277,30 @@ void AscendHardwareAssembler::GenerateMemcpyAsyncTrace(DataInventory &dataInvent
229 pidTidSet_.insert({formatPid, tid});277 pidTidSet_.insert({formatPid, tid});
230 std::shared_ptr<MemcpyAsyncEvent> event;278 std::shared_ptr<MemcpyAsyncEvent> event;
231 // 计算拷贝数据量和带宽279 // 计算拷贝数据量和带宽
232- if (showFlag) {280+ if (showFlag)
281+ {
233 TaskId taskId(data.streamId, data.batchId, data.taskId, data.contextId, data.deviceId);282 TaskId taskId(data.streamId, data.batchId, data.taskId, data.contextId, data.deviceId);
234 auto it = memcpyInfoDataMap.find(taskId);283 auto it = memcpyInfoDataMap.find(taskId);
235- if (it != memcpyInfoDataMap.end()) {284+ if (it != memcpyInfoDataMap.end())
285+ {
236 dataSize = it->second.dataSize;286 dataSize = it->second.dataSize;
237- memcpyDirection = it->second.memcpyOperation > VALID_MEMCPY_OPERATION ?287+ memcpyDirection = it->second.memcpyOperation > VALID_MEMCPY_OPERATION
238- OTHER_DIRECTION : MEMCPY_OPERATIONS[it->second.memcpyOperation];288+ ? OTHER_DIRECTION
239- } else {289+ : MEMCPY_OPERATIONS[it->second.memcpyOperation];
290+ }
291+ else
292+ {
240 ERROR("MEMCPY_ASYNC task lost memcpyInfo, connectionId is %", data.connectionId);293 ERROR("MEMCPY_ASYNC task lost memcpyInfo, connectionId is %", data.connectionId);
241 }294 }
242- if (!IsDoubleEqual(data.duration, 0.0) && data.duration > 0) {295+ if (!IsDoubleEqual(data.duration, 0.0) && data.duration > 0)
296+ {
243 bandwidth = static_cast<double>(dataSize) / data.duration; // GB/s, 全部按照1000计算297 bandwidth = static_cast<double>(dataSize) / data.duration; // GB/s, 全部按照1000计算
244 }298 }
245 }299 }
246 MAKE_SHARED_RETURN_VOID(event, MemcpyAsyncEvent, formatPid, tid, data.duration / NS_TO_US,300 MAKE_SHARED_RETURN_VOID(event, MemcpyAsyncEvent, formatPid, tid, data.duration / NS_TO_US,
247- DivideByPowersOfTenWithPrecision(data.timestamp), data.hostType,301+ DivideByPowersOfTenWithPrecision(data.timestamp), data.hostType, data.modelId,
248- data.modelId, data.streamId, data.taskId, data.batchId, data.contextId,302+ data.streamId, data.taskId, data.batchId, data.contextId, data.connectionId,
249- data.connectionId, data.deviceType, dataSize, bandwidth, memcpyDirection, showFlag);303+ data.deviceType, dataSize, bandwidth, memcpyDirection, showFlag);
250 res_.push_back(event);304 res_.push_back(event);
251 GenerateMemcpyAsyncConnectionTrace(data, formatPid);305 GenerateMemcpyAsyncConnectionTrace(data, formatPid);
252 }306 }
@@ -257,12 +311,13 @@ void AscendHardwareAssembler::GenerateTaskConnectionTrace(const AscendTaskData &
257 std::string connId;311 std::string connId;
258 std::string name;312 std::string name;
259 int tid;313 int tid;
260- if (opName_.find(id) != opName_.end() || aclEvent_.find(data.connectionId) != aclEvent_.end()) {314+ if (opName_.find(id) != opName_.end() || aclEvent_.find(data.connectionId) != aclEvent_.end())
315+ {
261 connId = ConnectionIdPool::GetConnectionId(data.connectionId, ConnectionCategory::GENERAL);316 connId = ConnectionIdPool::GetConnectionId(data.connectionId, ConnectionCategory::GENERAL);
262 name = HOST_TO_DEVICE + connId;317 name = HOST_TO_DEVICE + connId;
263 tid = static_cast<int>(GetPhysicStreamId(data.streamId));318 tid = static_cast<int>(GetPhysicStreamId(data.streamId));
264 std::shared_ptr<FlowEvent> end;319 std::shared_ptr<FlowEvent> end;
265- MAKE_SHARED_RETURN_VOID(end, FlowEvent, formatPid, tid, DivideByPowersOfTenWithPrecision(data.timestamp),320+ MAKE_SHARED_RETURN_VOID(end, FlowEvent, formatPid, tid, DivideByPowersOfTenWithPrecision(data.timestamp),
266 HOST_TO_DEVICE, connId, name, FLOW_END, FLOW_BP);321 HOST_TO_DEVICE, connId, name, FLOW_END, FLOW_BP);
267 res_.push_back(end);322 res_.push_back(end);
268 }323 }
@@ -274,43 +329,49 @@ void AscendHardwareAssembler::GenerateMemcpyAsyncConnectionTrace(const AscendTas
274 std::string name = HOST_TO_DEVICE + connId;329 std::string name = HOST_TO_DEVICE + connId;
275 int tid = static_cast<int>(GetPhysicStreamId(data.streamId));330 int tid = static_cast<int>(GetPhysicStreamId(data.streamId));
276 std::shared_ptr<FlowEvent> end;331 std::shared_ptr<FlowEvent> end;
277- MAKE_SHARED_RETURN_VOID(end, FlowEvent, formatPid, tid, DivideByPowersOfTenWithPrecision(data.timestamp),332+ MAKE_SHARED_RETURN_VOID(end, FlowEvent, formatPid, tid, DivideByPowersOfTenWithPrecision(data.timestamp),
278 HOST_TO_DEVICE, connId, name, FLOW_END, FLOW_BP);333 HOST_TO_DEVICE, connId, name, FLOW_END, FLOW_BP);
279 res_.push_back(end);334 res_.push_back(end);
280}335}
281 336 
282-uint8_t AscendHardwareAssembler::AssembleData(DataInventory& dataInventory, JsonWriter& ostream,337+uint8_t AscendHardwareAssembler::AssembleData(DataInventory &dataInventory, JsonWriter &ostream,
283- const std::string& profPath)338+ const std::string &profPath)
284{339{
285 auto taskData = dataInventory.GetPtr<std::vector<AscendTaskData>>();340 auto taskData = dataInventory.GetPtr<std::vector<AscendTaskData>>();
286 auto kfcTurnData = dataInventory.GetPtr<std::vector<KfcTurnData>>();341 auto kfcTurnData = dataInventory.GetPtr<std::vector<KfcTurnData>>();
287- if (taskData == nullptr && kfcTurnData == nullptr) {342+ if (taskData == nullptr && kfcTurnData == nullptr)
343+ {
288 WARN("Can't get task data from dataInventory");344 WARN("Can't get task data from dataInventory");
289 return DATA_NOT_EXIST;345 return DATA_NOT_EXIST;
290 }346 }
291 std::unordered_map<uint16_t, uint32_t> devicePid;347 std::unordered_map<uint16_t, uint32_t> devicePid;
292 auto layer = GetLayerInfo(PROCESS_TASK);348 auto layer = GetLayerInfo(PROCESS_TASK);
293- if (taskData != nullptr) {349+ if (taskData != nullptr)
350+ {
294 InitData(dataInventory, *taskData);351 InitData(dataInventory, *taskData);
295 GenerateTaskTrace(*taskData, profPath, layer, devicePid);352 GenerateTaskTrace(*taskData, profPath, layer, devicePid);
296- if (!memcpyAsyncDeviceTasks_.empty()) {353+ if (!memcpyAsyncDeviceTasks_.empty())
354+ {
297 GenerateMemcpyAsyncTrace(dataInventory, profPath, layer, devicePid);355 GenerateMemcpyAsyncTrace(dataInventory, profPath, layer, devicePid);
298 }356 }
299 }357 }
300- if (kfcTurnData != nullptr) {358+ if (kfcTurnData != nullptr)
359+ {
301 GenerateKfcTrace(*kfcTurnData, profPath, layer, devicePid);360 GenerateKfcTrace(*kfcTurnData, profPath, layer, devicePid);
302 }361 }
303 GenerateTaskMetaData(devicePid, layer, res_, pidTidSet_);362 GenerateTaskMetaData(devicePid, layer, res_, pidTidSet_);
304- if (res_.empty()) {363+ if (res_.empty())
364+ {
305 ERROR("Can't Generate any Ascend process data");365 ERROR("Can't Generate any Ascend process data");
306 return ASSEMBLE_FAILED;366 return ASSEMBLE_FAILED;
307 }367 }
308- for (const auto &node : res_) {368+ for (const auto &node : res_)
369+ {
309 node->DumpJson(ostream);370 node->DumpJson(ostream);
310 }371 }
311 // 为了让下一个写入的内容形成正确的JSON格式，需要补一个","372 // 为了让下一个写入的内容形成正确的JSON格式，需要补一个","
312 ostream << ",";373 ostream << ",";
313 return ASSEMBLE_SUCCESS;374 return ASSEMBLE_SUCCESS;
314}375}
315-}376+} // namespace Application
316-}377+} // namespace Analysis
Manalysis/csrc/application/timeline/ascend_hardware_assembler.h+82-32
@@ -19,27 +19,39 @@
19 19 
20#include <map>20#include <map>
21#include <unordered_set>21#include <unordered_set>
22+ 
22#include "analysis/csrc/application/timeline/json_assembler.h"23#include "analysis/csrc/application/timeline/json_assembler.h"
23#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/ascend_task_data.h"24#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/ascend_task_data.h"
24-#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/task_info_data.h"
25#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/kfc_turn_data.h"25#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/kfc_turn_data.h"
26+#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/task_info_data.h"
26#include "analysis/csrc/domain/valueobject/include/task_id.h"27#include "analysis/csrc/domain/valueobject/include/task_id.h"
27 28 
28-namespace Analysis {29+namespace Analysis
29-namespace Application {30+{
31+namespace Application
32+{
30 33 
31-class TaskTraceEvent : public DurationEvent {34+class TaskTraceEvent : public DurationEvent
32-public:35+{
36+ public:
33 TaskTraceEvent(uint32_t pid, int tid, double dur, const std::string &ts, const std::string &name, uint32_t modelId,37 TaskTraceEvent(uint32_t pid, int tid, double dur, const std::string &ts, const std::string &name, uint32_t modelId,
34 uint32_t streamId, uint32_t taskId, uint32_t batchId, uint32_t contextId, uint64_t connectionId,38 uint32_t streamId, uint32_t taskId, uint32_t batchId, uint32_t contextId, uint64_t connectionId,
35 const std::string taskType)39 const std::string taskType)
36- : DurationEvent(pid, tid, dur, ts, name), modelId_(modelId), streamId_(streamId), taskId_(taskId),40+ : DurationEvent(pid, tid, dur, ts, name),
37- batchId_(batchId), contextId_(contextId), connectionId_(connectionId), taskType_(taskType) {}41+ modelId_(modelId),
42+ streamId_(streamId),
43+ taskId_(taskId),
44+ batchId_(batchId),
45+ contextId_(contextId),
46+ connectionId_(connectionId),
47+ taskType_(taskType)
48+ {
49+ }
38 50 
39-protected:51+ protected:
40 void ProcessArgs(JsonWriter &ostream) override;52 void ProcessArgs(JsonWriter &ostream) override;
41 53 
42-private:54+ private:
43 uint32_t modelId_;55 uint32_t modelId_;
44 uint32_t streamId_;56 uint32_t streamId_;
45 uint32_t taskId_;57 uint32_t taskId_;
@@ -49,63 +61,101 @@ private:
49 std::string taskType_;61 std::string taskType_;
50};62};
51 63 
52-class MemcpyAsyncEvent : public TaskTraceEvent {64+class MemcpyAsyncEvent : public TaskTraceEvent
53-public:65+{
66+ public:
54 MemcpyAsyncEvent(uint32_t pid, int tid, double dur, const std::string &ts, const std::string &name,67 MemcpyAsyncEvent(uint32_t pid, int tid, double dur, const std::string &ts, const std::string &name,
55 uint32_t modelId, uint32_t streamId, uint32_t taskId, uint32_t batchId, uint32_t contextId,68 uint32_t modelId, uint32_t streamId, uint32_t taskId, uint32_t batchId, uint32_t contextId,
56 uint64_t connectionId, const std::string taskType, uint64_t dataSize, double bandwidth,69 uint64_t connectionId, const std::string taskType, uint64_t dataSize, double bandwidth,
57 std::string memcpyDirection, bool showFlag)70 std::string memcpyDirection, bool showFlag)
58 : TaskTraceEvent(pid, tid, dur, ts, name, modelId, streamId, taskId, batchId, contextId, connectionId,71 : TaskTraceEvent(pid, tid, dur, ts, name, modelId, streamId, taskId, batchId, contextId, connectionId,
59- taskType), dataSize_(dataSize), bandwidth_(bandwidth), memcpyDirection_(memcpyDirection),72+ taskType),
60- showFlag_(showFlag) {}73+ dataSize_(dataSize),
61-private:74+ bandwidth_(bandwidth),
75+ memcpyDirection_(memcpyDirection),
76+ showFlag_(showFlag)
77+ {
78+ }
79+ 
80+ private:
62 void ProcessArgs(JsonWriter &ostream) override;81 void ProcessArgs(JsonWriter &ostream) override;
63-private:82+ 
83+ private:
64 uint64_t dataSize_;84 uint64_t dataSize_;
65 double bandwidth_;85 double bandwidth_;
66 std::string memcpyDirection_;86 std::string memcpyDirection_;
67 bool showFlag_;87 bool showFlag_;
68};88};
69 89 
70-class KfcTurnTraceEvent : public DurationEvent {90+class SimtTaskEvent : public TaskTraceEvent
71-public:91+{
92+ public:
93+ SimtTaskEvent(uint32_t pid, int tid, double dur, const std::string &ts, const std::string &name, uint32_t modelId,
94+ uint32_t streamId, uint32_t taskId, uint32_t batchId, uint32_t contextId, uint64_t connectionId,
95+ const std::string taskType, std::string gridDim, std::string blockDim)
96+ : TaskTraceEvent(pid, tid, dur, ts, name, modelId, streamId, taskId, batchId, contextId, connectionId,
97+ taskType),
98+ gridDim_(std::move(gridDim)),
99+ blockDim_(std::move(blockDim))
100+ {
101+ }
102+ 
103+ private:
104+ void ProcessArgs(JsonWriter &ostream) override;
105+ 
106+ private:
107+ std::string gridDim_;
108+ std::string blockDim_;
109+};
110+ 
111+class KfcTurnTraceEvent : public DurationEvent
112+{
113+ public:
72 KfcTurnTraceEvent(int pid, int tid, double dur, const std::string &ts, const std::string &name, uint32_t streamId,114 KfcTurnTraceEvent(int pid, int tid, double dur, const std::string &ts, const std::string &name, uint32_t streamId,
73 uint32_t taskId)115 uint32_t taskId)
74- : DurationEvent(pid, tid, dur, ts, name), streamId_(streamId), taskId_(taskId) {}116+ : DurationEvent(pid, tid, dur, ts, name), streamId_(streamId), taskId_(taskId)
75-private:117+ {
118+ }
119+ 
120+ private:
76 void ProcessArgs(JsonWriter &ostream) override;121 void ProcessArgs(JsonWriter &ostream) override;
77-private:122+ 
123+ private:
78 uint32_t streamId_;124 uint32_t streamId_;
79 uint32_t taskId_;125 uint32_t taskId_;
80};126};
81 127 
82-class AscendHardwareAssembler : public JsonAssembler {128+class AscendHardwareAssembler : public JsonAssembler
83-public:129+{
130+ public:
84 AscendHardwareAssembler();131 AscendHardwareAssembler();
85-private:132+ 
86- uint8_t AssembleData(DataInventory& dataInventory, JsonWriter &ostream, const std::string &profPath) override;133+ private:
134+ uint8_t AssembleData(DataInventory &dataInventory, JsonWriter &ostream, const std::string &profPath) override;
87 void GenerateTaskTrace(const std::vector<AscendTaskData> &taskData, const std::string &profPath,135 void GenerateTaskTrace(const std::vector<AscendTaskData> &taskData, const std::string &profPath,
88 const LayerInfo &layer, std::unordered_map<uint16_t, uint32_t> &pidMap);136 const LayerInfo &layer, std::unordered_map<uint16_t, uint32_t> &pidMap);
89 void InitData(DataInventory &dataInventory, std::vector<AscendTaskData> &taskData);137 void InitData(DataInventory &dataInventory, std::vector<AscendTaskData> &taskData);
90- std::string GetOpName(const AscendTaskData& data);138+ std::string GetOpName(const AscendTaskData &data);
91- std::string GetTaskType(const AscendTaskData& data);139+ std::string GetTaskType(const AscendTaskData &data);
92 uint32_t GetPhysicStreamId(const uint32_t streamId);140 uint32_t GetPhysicStreamId(const uint32_t streamId);
93 void GenerateTaskConnectionTrace(const AscendTaskData &data, uint32_t formatPid, TaskId &id);141 void GenerateTaskConnectionTrace(const AscendTaskData &data, uint32_t formatPid, TaskId &id);
94- void GenerateKfcTrace(const std::vector<KfcTurnData>& kfcData, const std::string &profPath,142+ void GenerateKfcTrace(const std::vector<KfcTurnData> &kfcData, const std::string &profPath, const LayerInfo &layer,
95- const LayerInfo &layer, std::unordered_map<uint16_t, uint32_t> &pidMap);143+ std::unordered_map<uint16_t, uint32_t> &pidMap);
96 void GenerateMemcpyAsyncTrace(DataInventory &dataInventory, const std::string &profPath, const LayerInfo &layer,144 void GenerateMemcpyAsyncTrace(DataInventory &dataInventory, const std::string &profPath, const LayerInfo &layer,
97 std::unordered_map<uint16_t, uint32_t> &pidMap);145 std::unordered_map<uint16_t, uint32_t> &pidMap);
98 void GenerateMemcpyAsyncConnectionTrace(const AscendTaskData &data, uint32_t formatPid);146 void GenerateMemcpyAsyncConnectionTrace(const AscendTaskData &data, uint32_t formatPid);
99-private:147+ 
148+ private:
100 std::vector<std::shared_ptr<TraceEvent>> res_;149 std::vector<std::shared_ptr<TraceEvent>> res_;
101 std::shared_ptr<std::unordered_map<uint32_t, uint32_t>> logicStream_;150 std::shared_ptr<std::unordered_map<uint32_t, uint32_t>> logicStream_;
102 std::map<TaskId, std::string> opName_;151 std::map<TaskId, std::string> opName_;
103 std::map<TaskId, std::string> taskType_;152 std::map<TaskId, std::string> taskType_;
153+ std::map<TaskId, std::pair<std::string, std::string>> simtInfoMap_;
104 std::set<std::pair<uint32_t, int>> pidTidSet_;154 std::set<std::pair<uint32_t, int>> pidTidSet_;
105 std::set<TaskId> ffts_;155 std::set<TaskId> ffts_;
106 std::set<uint64_t> aclEvent_;156 std::set<uint64_t> aclEvent_;
107 std::vector<AscendTaskData> memcpyAsyncDeviceTasks_;157 std::vector<AscendTaskData> memcpyAsyncDeviceTasks_;
108};158};
109-}159+} // namespace Application
110-}160+} // namespace Analysis
111-#endif // ANALYSIS_APPLICATION_ASCEND_ASSEMBLER_H161+#endif // ANALYSIS_APPLICATION_ASCEND_ASSEMBLER_H
Manalysis/csrc/application/timeline/json_constant.h+10-6
@@ -17,11 +17,14 @@
17#ifndef ANALYSIS_APPLICATION_DELIVERABLES_CONSTANT_H17#ifndef ANALYSIS_APPLICATION_DELIVERABLES_CONSTANT_H
18#define ANALYSIS_APPLICATION_DELIVERABLES_CONSTANT_H18#define ANALYSIS_APPLICATION_DELIVERABLES_CONSTANT_H
19 19 
20-#include <string>
21#include <stdint.h>20#include <stdint.h>
22 21 
23-namespace Analysis {22+#include <string>
24-namespace Application {23+ 
24+namespace Analysis
25+{
26+namespace Application
27+{
25const uint8_t ASSEMBLE_FAILED = 0;28const uint8_t ASSEMBLE_FAILED = 0;
26const uint8_t ASSEMBLE_SUCCESS = 1;29const uint8_t ASSEMBLE_SUCCESS = 1;
27const uint8_t DATA_NOT_EXIST = 2;30const uint8_t DATA_NOT_EXIST = 2;
@@ -50,11 +53,12 @@ const std::string STEP_TRACE_FILE = "step_trace";
50const std::string MSPROF_TX_FILE = "msprof_tx";53const std::string MSPROF_TX_FILE = "msprof_tx";
51const std::string RECORD_EVENT = "aclrtRecordEvent";54const std::string RECORD_EVENT = "aclrtRecordEvent";
52const std::string MEMCPY_ASYNC = "MEMCPY_ASYNC";55const std::string MEMCPY_ASYNC = "MEMCPY_ASYNC";
56+const std::string KERNEL_SIMT_TASK_TYPE = "KERNEL_SIMT";
53/*57/*
54 * json格式要求多个对象使用[]包装，再每一层json后添加了","分割，最终会形成[{},true]的结果，因此需要写入内容的时候过滤掉58 * json格式要求多个对象使用[]包装，再每一层json后添加了","分割，最终会形成[{},true]的结果，因此需要写入内容的时候过滤掉
55 * [ true],共6位长度59 * [ true],共6位长度
56 */60 */
57const std::size_t FILE_CONTENT_SUFFIX = 6;61const std::size_t FILE_CONTENT_SUFFIX = 6;
58-}62+} // namespace Application
59-}63+} // namespace Analysis
60-#endif // ANALYSIS_APPLICATION_DELIVERABLES_CONSTANT_H64+#endif // ANALYSIS_APPLICATION_DELIVERABLES_CONSTANT_H
Manalysis/csrc/domain/data_process/ai_task/compute_task_info_processor.cpp+43-27
@@ -16,8 +16,10 @@
16 16 
17#include "analysis/csrc/domain/data_process/ai_task/compute_task_info_processor.h"17#include "analysis/csrc/domain/data_process/ai_task/compute_task_info_processor.h"
18 18 
19-namespace Analysis {19+namespace Analysis
20-namespace Domain {20+{
21+namespace Domain
22+{
21using namespace Analysis::Viewer::Database;23using namespace Analysis::Viewer::Database;
22ComputeTaskInfoProcessor::ComputeTaskInfoProcessor(const std::string &profPath) : DataProcessor(profPath) {}24ComputeTaskInfoProcessor::ComputeTaskInfoProcessor(const std::string &profPath) : DataProcessor(profPath) {}
23 25 
@@ -25,69 +27,83 @@ bool ComputeTaskInfoProcessor::Process(DataInventory &dataInventory)
25{27{
26 DBInfo geInfoDB("ge_info.db", "TaskInfo");28 DBInfo geInfoDB("ge_info.db", "TaskInfo");
27 std::string dbPath = Utils::File::PathJoin({profPath_, HOST, SQLITE, geInfoDB.dbName});29 std::string dbPath = Utils::File::PathJoin({profPath_, HOST, SQLITE, geInfoDB.dbName});
28- if (!geInfoDB.ConstructDBRunner(dbPath)) {30+ if (!geInfoDB.ConstructDBRunner(dbPath))
31+ {
29 return false;32 return false;
30 }33 }
31 // 并不是所有场景都有ge info数据34 // 并不是所有场景都有ge info数据
32 auto status = CheckPathAndTable(dbPath, geInfoDB);35 auto status = CheckPathAndTable(dbPath, geInfoDB);
33- if (status == CHECK_FAILED) {36+ if (status == CHECK_FAILED)
37+ {
34 return false;38 return false;
35- } else if (status == NOT_EXIST) {39+ }
40+ else if (status == NOT_EXIST)
41+ {
36 return true;42 return true;
37 }43 }
38 auto hashMap = dataInventory.GetPtr<std::unordered_map<std::string, std::string>>();44 auto hashMap = dataInventory.GetPtr<std::unordered_map<std::string, std::string>>();
39- if (hashMap == nullptr) {45+ if (hashMap == nullptr)
46+ {
40 ERROR("Can't get hash data.");47 ERROR("Can't get hash data.");
41 return false;48 return false;
42 }49 }
43 auto formatData = LoadData(geInfoDB, dbPath, *hashMap);50 auto formatData = LoadData(geInfoDB, dbPath, *hashMap);
44- if (formatData.empty()) {51+ if (formatData.empty())
52+ {
45 ERROR("TaskInfo format data is empty. DBPath is %", dbPath);53 ERROR("TaskInfo format data is empty. DBPath is %", dbPath);
46 return false;54 return false;
47 }55 }
48- if (!SaveToDataInventory<TaskInfoData>(std::move(formatData), dataInventory,56+ if (!SaveToDataInventory<TaskInfoData>(std::move(formatData), dataInventory, PROCESSOR_NAME_COMPUTE_TASK_INFO))
49- PROCESSOR_NAME_COMPUTE_TASK_INFO)) {57+ {
50 ERROR("Save data failed, %.", PROCESSOR_NAME_COMPUTE_TASK_INFO);58 ERROR("Save data failed, %.", PROCESSOR_NAME_COMPUTE_TASK_INFO);
51 return false;59 return false;
52 }60 }
53 return true;61 return true;
54}62}
55 63 
56-std::vector<TaskInfoData> ComputeTaskInfoProcessor::LoadData(64+std::vector<TaskInfoData> ComputeTaskInfoProcessor::LoadData(const DBInfo &taskInfoDB, const std::string &dbPath,
57- const DBInfo &taskInfoDB, const std::string &dbPath, std::unordered_map<std::string, std::string>& hashMap)65+ std::unordered_map<std::string, std::string> &hashMap)
58{66{
59 std::vector<TaskInfoFormat> oriData;67 std::vector<TaskInfoFormat> oriData;
60 std::vector<TaskInfoData> res;68 std::vector<TaskInfoData> res;
61- if (taskInfoDB.dbRunner == nullptr) {69+ if (taskInfoDB.dbRunner == nullptr)
70+ {
62 ERROR("Create % connection failed.", dbPath);71 ERROR("Create % connection failed.", dbPath);
63 return res;72 return res;
64 }73 }
65- std::string sql{"SELECT hashid, model_id, op_name, stream_id, task_id, block_num, mix_block_num, task_type, "74+ std::string sql{
66- "op_type, op_flag, batch_id, IFNULL(input_formats, 'NULL'), IFNULL(input_data_types, 'NULL'), "75+ "SELECT hashid, model_id, op_name, stream_id, task_id, block_num, mix_block_num, task_type, "
67- "IFNULL(input_shapes, 'NULL'), IFNULL(output_formats, 'NULL'), IFNULL(output_data_types, 'NULL'), "76+ "op_type, op_flag, batch_id, IFNULL(input_formats, 'NULL'), IFNULL(input_data_types, 'NULL'), "
68- "IFNULL(output_shapes, 'NULL'), device_id, context_id, (case when op_state is 'N/A' then 'N/A' "77+ "IFNULL(input_shapes, 'NULL'), IFNULL(output_formats, 'NULL'), IFNULL(output_data_types, 'NULL'), "
69- "when op_state is '1' then 'dynamic' when op_state is '0' then 'static' end) FROM " +78+ "IFNULL(output_shapes, 'NULL'), device_id, context_id, (case when op_state is 'N/A' then 'N/A' "
70- taskInfoDB.tableName};79+ "when op_state is '1' then 'dynamic' when op_state is '0' then 'static' end), "
71- if (!taskInfoDB.dbRunner->QueryData(sql, oriData)) {80+ "IFNULL(grid_dim, 'N/A'), IFNULL(block_dim, 'N/A') FROM " +
81+ taskInfoDB.tableName};
82+ if (!taskInfoDB.dbRunner->QueryData(sql, oriData))
83+ {
72 ERROR("Failed to obtain data from the % table.", taskInfoDB.tableName);84 ERROR("Failed to obtain data from the % table.", taskInfoDB.tableName);
73 }85 }
74- if (oriData.empty()) {86+ if (oriData.empty())
87+ {
75 ERROR("TaskInfo original data is empty, DBPath is %", dbPath);88 ERROR("TaskInfo original data is empty, DBPath is %", dbPath);
76 return res;89 return res;
77 }90 }
78- if (!Utils::Reserve(res, oriData.size())) {91+ if (!Utils::Reserve(res, oriData.size()))
92+ {
79 ERROR("Reserve for TaskInfo data failed");93 ERROR("Reserve for TaskInfo data failed");
80 return res;94 return res;
81 }95 }
82 TaskInfoData data;96 TaskInfoData data;
83 ShapeInfo info;97 ShapeInfo info;
84 std::string hashId;98 std::string hashId;
85- for (auto& node : oriData) {99+ for (auto &node : oriData)
100+ {
86 std::tie(hashId, data.modelId, data.opName, data.streamId, data.taskId, data.blockNum, data.mixBlockNum,101 std::tie(hashId, data.modelId, data.opName, data.streamId, data.taskId, data.blockNum, data.mixBlockNum,
87 data.taskType, data.opType, data.opFlag, data.batchId, info.inputFormats, info.inputDataTypes,102 data.taskType, data.opType, data.opFlag, data.batchId, info.inputFormats, info.inputDataTypes,
88- info.inputShapes, info.outputFormats, info.outputDataTypes, info.outputShapes,103+ info.inputShapes, info.outputFormats, info.outputDataTypes, info.outputShapes, data.deviceId,
89- data.deviceId, data.contextId, data.opState) = node;104+ data.contextId, data.opState, data.gridDim, data.blockDim) = node;
90- if (hashMap.find(hashId) != hashMap.end()) {105+ if (hashMap.find(hashId) != hashMap.end())
106+ {
91 data.hashId = hashMap[hashId];107 data.hashId = hashMap[hashId];
92 }108 }
93 data = info;109 data = info;
@@ -95,5 +111,5 @@ std::vector<TaskInfoData> ComputeTaskInfoProcessor::LoadData(
95 }111 }
96 return res;112 return res;
97}113}
98-}114+} // namespace Domain
99-}115+} // namespace Analysis
Manalysis/csrc/domain/data_process/ai_task/compute_task_info_processor.h+17-13
@@ -19,25 +19,29 @@
19#include "analysis/csrc/domain/data_process/data_processor.h"19#include "analysis/csrc/domain/data_process/data_processor.h"
20#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/task_info_data.h"20#include "analysis/csrc/domain/entities/viewer_data/ai_task/include/task_info_data.h"
21 21 
22-namespace Analysis {22+namespace Analysis
23-namespace Domain {23+{
24+namespace Domain
25+{
24// hashid, model_id, op_name, stream_id, task_id, block_num, mix_block_num, task_type, op_type,26// hashid, model_id, op_name, stream_id, task_id, block_num, mix_block_num, task_type, op_type,
25// op_flag, batch_id, input_formats, input_data_types, input_shapes, output_formats,27// op_flag, batch_id, input_formats, input_data_types, input_shapes, output_formats,
26-// output_data_types, output_shapes, device_id, context_id, op_state28+// output_data_types, output_shapes, device_id, context_id, op_state, grid_dim, block_dim
27-using TaskInfoFormat = std::tuple<std::string, uint32_t, std::string, uint32_t, uint32_t, uint32_t, uint32_t,29+using TaskInfoFormat =
28- std::string, std::string, std::string, uint32_t, std::string, std::string,30+ std::tuple<std::string, uint32_t, std::string, uint32_t, uint32_t, uint32_t, uint32_t, std::string, std::string,
29- std::string, std::string, std::string, std::string, uint16_t, uint32_t, std::string>;31+ std::string, uint32_t, std::string, std::string, std::string, std::string, std::string, std::string,
30-class ComputeTaskInfoProcessor : public DataProcessor {32+ uint16_t, uint32_t, std::string, std::string, std::string>;
31-public:33+class ComputeTaskInfoProcessor : public DataProcessor
34+{
35+ public:
32 ComputeTaskInfoProcessor() = default;36 ComputeTaskInfoProcessor() = default;
33 explicit ComputeTaskInfoProcessor(const std::string &profPath);37 explicit ComputeTaskInfoProcessor(const std::string &profPath);
34 38 
35-private:39+ private:
36 bool Process(DataInventory &dataInventory) override;40 bool Process(DataInventory &dataInventory) override;
37 std::vector<TaskInfoData> LoadData(const DBInfo &taskInfoDB, const std::string &dbPath,41 std::vector<TaskInfoData> LoadData(const DBInfo &taskInfoDB, const std::string &dbPath,
38- std::unordered_map<std::string, std::string>& hashMap);42+ std::unordered_map<std::string, std::string> &hashMap);
39};43};
40-}44+} // namespace Domain
41-}45+} // namespace Analysis
42 46 
43-#endif // ANALYSIS_DOMAIN_COMPUTE_TASK_INFO_PROCESSOR_H47+#endif // ANALYSIS_DOMAIN_COMPUTE_TASK_INFO_PROCESSOR_H
Manalysis/csrc/domain/entities/hal/include/ascend_obj.h+95-44
@@ -19,18 +19,21 @@
19 19 
20#include <cstdint>20#include <cstdint>
21#include <memory>21#include <memory>
22-#include <utility>
23#include <string>22#include <string>
23+#include <utility>
24 24 
25#include "analysis/csrc/infrastructure/utils/prof_common.h"25#include "analysis/csrc/infrastructure/utils/prof_common.h"
26 26 
27-namespace Analysis {27+namespace Analysis
28-namespace Domain {28+{
29+namespace Domain
30+{
29 31 
30#define DEFAULT_CONTEXT_ID 0xffffffff32#define DEFAULT_CONTEXT_ID 0xffffffff
31 33 
32// 算子类型枚举34// 算子类型枚举
33-enum class OpType {35+enum class OpType
36+{
34 OPTYPE_HCCL_BIG = 0,37 OPTYPE_HCCL_BIG = 0,
35 OPTYPE_HCCL_SMALL,38 OPTYPE_HCCL_SMALL,
36 OPTYPE_COMPUTE,39 OPTYPE_COMPUTE,
@@ -40,26 +43,31 @@ enum class OpType {
40};43};
41 44 
42// 计算算子描述信息45// 计算算子描述信息
43-struct OpDesc {46+struct OpDesc
47+{
44 std::shared_ptr<MsprofCompactInfo> nodeDesc = nullptr;48 std::shared_ptr<MsprofCompactInfo> nodeDesc = nullptr;
45 std::shared_ptr<MsprofCompactInfo> nodeAttr = nullptr;49 std::shared_ptr<MsprofCompactInfo> nodeAttr = nullptr;
50+ std::shared_ptr<MsprofCompactInfo> runtimeTrackDesc = nullptr;
46 std::shared_ptr<ConcatTensorInfo> tensorDesc = nullptr;51 std::shared_ptr<ConcatTensorInfo> tensorDesc = nullptr;
47 std::shared_ptr<MsprofAdditionalInfo> ctxId = nullptr;52 std::shared_ptr<MsprofAdditionalInfo> ctxId = nullptr;
48};53};
49 54 
50// 通信小算子描述信息55// 通信小算子描述信息
51-struct HcclSmallOpDesc {56+struct HcclSmallOpDesc
57+{
52 uint32_t ctxId = DEFAULT_CONTEXT_ID;58 uint32_t ctxId = DEFAULT_CONTEXT_ID;
53 uint8_t isMaster = 0; // 1代表master59 uint8_t isMaster = 0; // 1代表master
54 std::shared_ptr<MsprofAdditionalInfo> hcclInfo = nullptr;60 std::shared_ptr<MsprofAdditionalInfo> hcclInfo = nullptr;
55 61 
56 HcclSmallOpDesc(uint32_t ctxId, uint32_t isMaster, const std::shared_ptr<MsprofAdditionalInfo> &hcclInfo)62 HcclSmallOpDesc(uint32_t ctxId, uint32_t isMaster, const std::shared_ptr<MsprofAdditionalInfo> &hcclInfo)
57 : ctxId(ctxId), isMaster(isMaster), hcclInfo(hcclInfo)63 : ctxId(ctxId), isMaster(isMaster), hcclInfo(hcclInfo)
58- {}64+ {
65+ }
59};66};
60 67 
61// 通信大算子描述信息68// 通信大算子描述信息
62-struct HcclBigOpDesc {69+struct HcclBigOpDesc
70+{
63 uint64_t beginTime = 0;71 uint64_t beginTime = 0;
64 uint64_t endTime = 0;72 uint64_t endTime = 0;
65 uint16_t deviceId = 0;73 uint16_t deviceId = 0;
@@ -74,57 +82,76 @@ struct HcclBigOpDesc {
74 HcclBigOpDesc(uint64_t begin, uint64_t end, uint16_t deviceId, uint64_t modelId, int32_t indexId,82 HcclBigOpDesc(uint64_t begin, uint64_t end, uint16_t deviceId, uint64_t modelId, int32_t indexId,
75 int64_t connectionId, uint32_t threadId, const std::shared_ptr<MsprofCompactInfo> &node,83 int64_t connectionId, uint32_t threadId, const std::shared_ptr<MsprofCompactInfo> &node,
76 const std::shared_ptr<MsprofCompactInfo> &hcclOpDesc, int64_t kfcConnectionId)84 const std::shared_ptr<MsprofCompactInfo> &hcclOpDesc, int64_t kfcConnectionId)
77- : beginTime(begin), endTime(end), deviceId(deviceId), modelId(modelId), indexId(indexId),85+ : beginTime(begin),
78- connectionId(connectionId), thread_id(threadId), nodeDesc(node), opInfoDesc(hcclOpDesc),86+ endTime(end),
87+ deviceId(deviceId),
88+ modelId(modelId),
89+ indexId(indexId),
90+ connectionId(connectionId),
91+ thread_id(threadId),
92+ nodeDesc(node),
93+ opInfoDesc(hcclOpDesc),
79 kfcConnectionId(kfcConnectionId)94 kfcConnectionId(kfcConnectionId)
80- {}95+ {
96+ }
81};97};
82 98 
83// 算子统一对外结构体99// 算子统一对外结构体
84-struct Operator {100+struct Operator
101+{
85 Operator(std::shared_ptr<OpDesc> desc, const uint64_t &name, const OpType &type)102 Operator(std::shared_ptr<OpDesc> desc, const uint64_t &name, const OpType &type)
86 : opDesc(std::move(desc)), name(name), type(type)103 : opDesc(std::move(desc)), name(name), type(type)
87- {}104+ {
105+ }
88 106 
89 Operator(std::shared_ptr<HcclSmallOpDesc> desc, const uint64_t &name, const OpType &type)107 Operator(std::shared_ptr<HcclSmallOpDesc> desc, const uint64_t &name, const OpType &type)
90 : hcclSmallOpDesc(std::move(desc)), name(name), type(type)108 : hcclSmallOpDesc(std::move(desc)), name(name), type(type)
91- {}109+ {
110+ }
92 111 
93 Operator(std::shared_ptr<HcclBigOpDesc> desc, const uint64_t &name, const OpType &type)112 Operator(std::shared_ptr<HcclBigOpDesc> desc, const uint64_t &name, const OpType &type)
94 : hcclBigOpDesc(std::move(desc)), name(name), type(type)113 : hcclBigOpDesc(std::move(desc)), name(name), type(type)
95- {}114+ {
115+ }
96 116 
97- union {117+ union
118+ {
98 std::shared_ptr<OpDesc> opDesc;119 std::shared_ptr<OpDesc> opDesc;
99 std::shared_ptr<HcclSmallOpDesc> hcclSmallOpDesc;120 std::shared_ptr<HcclSmallOpDesc> hcclSmallOpDesc;
100 std::shared_ptr<HcclBigOpDesc> hcclBigOpDesc;121 std::shared_ptr<HcclBigOpDesc> hcclBigOpDesc;
101 };122 };
102 123 
103- uint64_t name = 0; // aka item_id124+ uint64_t name = 0; // aka item_id
104 OpType type = OpType::OPTYPE_INVALID;125 OpType type = OpType::OPTYPE_INVALID;
105 126 
106 ~Operator()127 ~Operator()
107 {128 {
108- if (opDesc) {129+ if (opDesc)
130+ {
109 opDesc.~shared_ptr();131 opDesc.~shared_ptr();
110- } else if (hcclSmallOpDesc) {132+ }
133+ else if (hcclSmallOpDesc)
134+ {
111 hcclSmallOpDesc.~shared_ptr();135 hcclSmallOpDesc.~shared_ptr();
112- } else if (hcclBigOpDesc) {136+ }
137+ else if (hcclBigOpDesc)
138+ {
113 hcclBigOpDesc.~shared_ptr();139 hcclBigOpDesc.~shared_ptr();
114 }140 }
115 }141 }
116};142};
117 143 
118// 用于存储分析树结果，DBDumper将此信息落盘144// 用于存储分析树结果，DBDumper将此信息落盘
119-struct HostTask {145+struct HostTask
120- uint32_t taskId = 0; // 和采集侧的数据类型不一致，采集侧的高低16位会被分别用作batchId和TaskId146+{
147+ uint32_t taskId = 0; // 和采集侧的数据类型不一致，采集侧的高低16位会被分别用作batchId和TaskId
121 uint16_t batchId = 0;148 uint16_t batchId = 0;
122 uint16_t deviceId = 0;149 uint16_t deviceId = 0;
123 int32_t requestId = 0;150 int32_t requestId = 0;
124 uint32_t thread_id = 0;151 uint32_t thread_id = 0;
125 uint32_t streamId = 0;152 uint32_t streamId = 0;
126 uint32_t contextId = 0;153 uint32_t contextId = 0;
127- int64_t connection_id = 0; // -1 表示该任务无node直连154+ int64_t connection_id = 0; // -1 表示该任务无node直连
128 uint64_t modelId = 0;155 uint64_t modelId = 0;
129 uint64_t taskType = 0;156 uint64_t taskType = 0;
130 uint64_t timeStamp = 0;157 uint64_t timeStamp = 0;
@@ -135,16 +162,19 @@ struct HostTask {
135};162};
136 163 
137// 存储GeFusionOpInfo表164// 存储GeFusionOpInfo表
138-struct GeFusionOpInfo {165+struct GeFusionOpInfo
166+{
139 uint64_t modelId;167 uint64_t modelId;
140 std::shared_ptr<ProfFusionOpInfo> fusionOpInfo = nullptr;168 std::shared_ptr<ProfFusionOpInfo> fusionOpInfo = nullptr;
141 GeFusionOpInfo(uint64_t modelId, const std::shared_ptr<ProfFusionOpInfo> &fusionOp)169 GeFusionOpInfo(uint64_t modelId, const std::shared_ptr<ProfFusionOpInfo> &fusionOp)
142 : modelId(modelId), fusionOpInfo(fusionOp)170 : modelId(modelId), fusionOpInfo(fusionOp)
143- {}171+ {
172+ }
144};173};
145 174 
146// 用于存储分析树结果，DBDumper将此信息落盘175// 用于存储分析树结果，DBDumper将此信息落盘
147-struct RuntimeOpInfo {176+struct RuntimeOpInfo
177+{
148 bool isValid = false;178 bool isValid = false;
149 uint16_t deviceId = 0;179 uint16_t deviceId = 0;
150 uint16_t taskId = 0;180 uint16_t taskId = 0;
@@ -168,22 +198,37 @@ struct RuntimeOpInfo {
168 198 
169 RuntimeOpInfo() = default;199 RuntimeOpInfo() = default;
170 RuntimeOpInfo(uint16_t deviceId, uint16_t taskId, uint16_t blockNum, uint16_t mixBlockNum, uint16_t opFlag,200 RuntimeOpInfo(uint16_t deviceId, uint16_t taskId, uint16_t blockNum, uint16_t mixBlockNum, uint16_t opFlag,
171- uint16_t tensorNum, uint32_t streamId, uint64_t modelId, std::string taskType,201+ uint16_t tensorNum, uint32_t streamId, uint64_t modelId, std::string taskType, std::string opType,
172- std::string opType, std::string opName, std::string hashId, std::string isDynamic,202+ std::string opName, std::string hashId, std::string isDynamic, std::string inputFormats,
173- std::string inputFormats, std::string inputDataTypes, std::string inputShapes,203+ std::string inputDataTypes, std::string inputShapes, std::string outputFormats,
174- std::string outputFormats, std::string outputDataTypes, std::string outputShapes)204+ std::string outputDataTypes, std::string outputShapes)
175- : deviceId(deviceId), taskId(taskId), blockNum(blockNum), mixBlockNum(mixBlockNum), opFlag(opFlag),205+ : deviceId(deviceId),
176- tensorNum(tensorNum), streamId(streamId), modelId(modelId), taskType(std::move(taskType)),206+ taskId(taskId),
177- opType(std::move(opType)), opName(std::move(opName)), hashId(std::move(hashId)),207+ blockNum(blockNum),
208+ mixBlockNum(mixBlockNum),
209+ opFlag(opFlag),
210+ tensorNum(tensorNum),
211+ streamId(streamId),
212+ modelId(modelId),
213+ taskType(std::move(taskType)),
214+ opType(std::move(opType)),
215+ opName(std::move(opName)),
216+ hashId(std::move(hashId)),
178 isDynamic(std::move(isDynamic)),217 isDynamic(std::move(isDynamic)),
179- inputFormats(std::move(inputFormats)), inputDataTypes(std::move(inputDataTypes)),218+ inputFormats(std::move(inputFormats)),
180- inputShapes(std::move(inputShapes)), outputFormats(std::move(outputFormats)),219+ inputDataTypes(std::move(inputDataTypes)),
181- outputDataTypes(std::move(outputDataTypes)), outputShapes(std::move(outputShapes)), isValid(true)220+ inputShapes(std::move(inputShapes)),
182- {}221+ outputFormats(std::move(outputFormats)),
222+ outputDataTypes(std::move(outputDataTypes)),
223+ outputShapes(std::move(outputShapes)),
224+ isValid(true)
225+ {
226+ }
183};227};
184 228 
185// 存储CaptureStreamInfo表229// 存储CaptureStreamInfo表
186-struct CaptureStreamInfo {230+struct CaptureStreamInfo
231+{
187 uint64_t modelId = UINT32_MAX;232 uint64_t modelId = UINT32_MAX;
188 uint64_t timeStamp = 0;233 uint64_t timeStamp = 0;
189 uint32_t streamId = 0;234 uint32_t streamId = 0;
@@ -195,11 +240,17 @@ struct CaptureStreamInfo {
195 CaptureStreamInfo() = default;240 CaptureStreamInfo() = default;
196 CaptureStreamInfo(uint64_t modelId, uint64_t timeStamp, uint32_t streamId, uint16_t originalStreamId,241 CaptureStreamInfo(uint64_t modelId, uint64_t timeStamp, uint32_t streamId, uint16_t originalStreamId,
197 uint16_t deviceId, uint16_t batchId, uint16_t captureStatus)242 uint16_t deviceId, uint16_t batchId, uint16_t captureStatus)
198- : modelId(modelId), timeStamp(timeStamp), streamId(streamId), originalStreamId(originalStreamId),243+ : modelId(modelId),
199- deviceId(deviceId), batchId(batchId), captureStatus(captureStatus)244+ timeStamp(timeStamp),
200- {}245+ streamId(streamId),
246+ originalStreamId(originalStreamId),
247+ deviceId(deviceId),
248+ batchId(batchId),
249+ captureStatus(captureStatus)
250+ {
251+ }
201};252};
202 253 
203-} // namespace Domain254+} // namespace Domain
204-} // namespace Analysis255+} // namespace Analysis
205-#endif // ANALYSIS_ENTITIES_ASCEND_OBJ_H256+#endif // ANALYSIS_ENTITIES_ASCEND_OBJ_H
Manalysis/csrc/domain/entities/viewer_data/ai_task/include/task_info_data.h+19-11
@@ -16,16 +16,20 @@
16#ifndef ANALYSIS_DOMAIN_TASK_INFO_DATA_H16#ifndef ANALYSIS_DOMAIN_TASK_INFO_DATA_H
17#define ANALYSIS_DOMAIN_TASK_INFO_DATA_H17#define ANALYSIS_DOMAIN_TASK_INFO_DATA_H
18 18 
19-#include <string>
20#include <stdint.h>19#include <stdint.h>
21 20 
22-namespace Analysis {21+#include <string>
23-namespace Domain {22+ 
23+namespace Analysis
24+{
25+namespace Domain
26+{
24const uint8_t DEFAULT_MULTIPLE_SIZE = 2;27const uint8_t DEFAULT_MULTIPLE_SIZE = 2;
25const std::string SINGLE_OPERATOR = "\"";28const std::string SINGLE_OPERATOR = "\"";
26const std::string CSV_OPERATOR = R"(""")";29const std::string CSV_OPERATOR = R"(""")";
27 30 
28-struct ShapeInfo {31+struct ShapeInfo
32+{
29 std::string inputFormats;33 std::string inputFormats;
30 std::string inputDataTypes;34 std::string inputDataTypes;
31 std::string inputShapes;35 std::string inputShapes;
@@ -34,7 +38,8 @@ struct ShapeInfo {
34 std::string outputShapes;38 std::string outputShapes;
35};39};
36 40 
37-struct TaskInfoData {41+struct TaskInfoData
42+{
38 uint16_t deviceId = UINT16_MAX;43 uint16_t deviceId = UINT16_MAX;
39 uint32_t modelId = UINT32_MAX;44 uint32_t modelId = UINT32_MAX;
40 uint32_t streamId = UINT32_MAX;45 uint32_t streamId = UINT32_MAX;
@@ -55,8 +60,10 @@ struct TaskInfoData {
55 std::string outputFormats;60 std::string outputFormats;
56 std::string outputDataTypes;61 std::string outputDataTypes;
57 std::string outputShapes;62 std::string outputShapes;
63+ std::string gridDim;
64+ std::string blockDim;
58 65 
59- TaskInfoData& operator=(const ShapeInfo &shapeInfo)66+ TaskInfoData &operator=(const ShapeInfo &shapeInfo)
60 {67 {
61 inputFormats = escapeQuotes(shapeInfo.inputFormats);68 inputFormats = escapeQuotes(shapeInfo.inputFormats);
62 inputDataTypes = escapeQuotes(shapeInfo.inputDataTypes);69 inputDataTypes = escapeQuotes(shapeInfo.inputDataTypes);
@@ -67,19 +74,20 @@ struct TaskInfoData {
67 return *this;74 return *this;
68 }75 }
69 76 
70-private:77+ private:
71 std::string escapeQuotes(const std::string &input)78 std::string escapeQuotes(const std::string &input)
72 {79 {
73 std::string res = input;80 std::string res = input;
74 res.reserve(input.size() * DEFAULT_MULTIPLE_SIZE);81 res.reserve(input.size() * DEFAULT_MULTIPLE_SIZE);
75 std::string::size_type pos = 0;82 std::string::size_type pos = 0;
76- while ((pos = res.find(SINGLE_OPERATOR, pos)) != std::string::npos) {83+ while ((pos = res.find(SINGLE_OPERATOR, pos)) != std::string::npos)
84+ {
77 res.replace(pos, SINGLE_OPERATOR.length(), CSV_OPERATOR);85 res.replace(pos, SINGLE_OPERATOR.length(), CSV_OPERATOR);
78 pos += CSV_OPERATOR.length();86 pos += CSV_OPERATOR.length();
79 }87 }
80 return res;88 return res;
81 }89 }
82};90};
83-}91+} // namespace Domain
84-}92+} // namespace Analysis
85-#endif // ANALYSIS_DOMAIN_TASK_INFO_DATA_H93+#endif // ANALYSIS_DOMAIN_TASK_INFO_DATA_H
Manalysis/csrc/domain/services/association/cann/tree_analyzer.cpp+273-201
@@ -15,21 +15,26 @@
15 * -------------------------------------------------------------------------*/15 * -------------------------------------------------------------------------*/
16#include "analysis/csrc/domain/services/association/cann/include/tree_analyzer.h"16#include "analysis/csrc/domain/services/association/cann/include/tree_analyzer.h"
17 17 
18-#include <string>
19#include <memory>18#include <memory>
20#include <set>19#include <set>
20+#include <string>
21+ 
21#include "analysis/csrc/domain/services/parser/host/cann/hash_data.h"22#include "analysis/csrc/domain/services/parser/host/cann/hash_data.h"
22#include "analysis/csrc/domain/services/parser/host/cann/rt_add_info_center.h"23#include "analysis/csrc/domain/services/parser/host/cann/rt_add_info_center.h"
23 24 
24-namespace Analysis {25+namespace Analysis
25-namespace Domain {26+{
26-namespace Cann {27+namespace Domain
28+{
29+namespace Cann
30+{
27 31 
28using namespace Analysis::Domain;32using namespace Analysis::Domain;
29using namespace Analysis::Utils;33using namespace Analysis::Utils;
30using namespace Analysis::Domain::Host::Cann;34using namespace Analysis::Domain::Host::Cann;
31 35 
32-namespace {36+namespace
37+{
33const uint32_t VALID_CTXID_NUM = 2;38const uint32_t VALID_CTXID_NUM = 2;
34const int64_t INVALID_VALUE = -1;39const int64_t INVALID_VALUE = -1;
35const uint16_t TWO_BYTES = 16;40const uint16_t TWO_BYTES = 16;
@@ -51,64 +56,50 @@ const std::string KERNEL_AI_VECTOR_CORE_TASK_TYPE = "KERNEL_AIVEC";
51const std::string KERNEL_MIX_AIC_TASK_TYPE = "KERNEL_MIX_AIC";56const std::string KERNEL_MIX_AIC_TASK_TYPE = "KERNEL_MIX_AIC";
52const std::string KERNEL_MIX_AIV_TASK_TYPE = "KERNEL_MIX_AIV";57const std::string KERNEL_MIX_AIV_TASK_TYPE = "KERNEL_MIX_AIV";
53const std::string KERNEL_AI_CPU_TASK_TYPE = "KERNEL_AICPU";58const std::string KERNEL_AI_CPU_TASK_TYPE = "KERNEL_AICPU";
54-const std::set<std::string> CONTEXT_ID_WHITE_LIST = {KERNEL_AI_CORE_TASK_TYPE, KERNEL_AI_VECTOR_CORE_TASK_TYPE,59+const std::string KERNEL_SIMT_TASK_TYPE = "KERNEL_SIMT";
55- KERNEL_FFTS_PLUS_TASK_TYPE,60+const std::set<std::string> CONTEXT_ID_WHITE_LIST = {KERNEL_AI_CORE_TASK_TYPE, KERNEL_AI_VECTOR_CORE_TASK_TYPE,
56- KERNEL_MIX_AIC_TASK_TYPE, KERNEL_MIX_AIV_TASK_TYPE};61+ KERNEL_FFTS_PLUS_TASK_TYPE, KERNEL_SIMT_TASK_TYPE,
62+ KERNEL_MIX_AIC_TASK_TYPE, KERNEL_MIX_AIV_TASK_TYPE};
57const std::set<std::string> KERNEL_COMPUTE_WHITE_LIST = {KERNEL_AI_CORE_TASK_TYPE, KERNEL_AI_VECTOR_CORE_TASK_TYPE,63const std::set<std::string> KERNEL_COMPUTE_WHITE_LIST = {KERNEL_AI_CORE_TASK_TYPE, KERNEL_AI_VECTOR_CORE_TASK_TYPE,
58- KERNEL_AI_CPU_TASK_TYPE,64+ KERNEL_AI_CPU_TASK_TYPE, KERNEL_SIMT_TASK_TYPE,
59 KERNEL_MIX_AIC_TASK_TYPE, KERNEL_MIX_AIV_TASK_TYPE};65 KERNEL_MIX_AIC_TASK_TYPE, KERNEL_MIX_AIV_TASK_TYPE};
60 66 
61uint64_t GetModelId(const std::shared_ptr<MsprofApi> &api, uint16_t deviceId, uint32_t streamId, uint16_t batchId,67uint64_t GetModelId(const std::shared_ptr<MsprofApi> &api, uint16_t deviceId, uint32_t streamId, uint16_t batchId,
62 uint64_t timestamp)68 uint64_t timestamp)
63{69{
64- return api != nullptr ? api->itemId : RTAddInfoCenter::GetInstance().GetModelId(deviceId, streamId, batchId,70+ return api != nullptr ? api->itemId
65- timestamp);71+ : RTAddInfoCenter::GetInstance().GetModelId(deviceId, streamId, batchId, timestamp);
66}72}
67 73 
68-}74+} // namespace
69 75 
70-void TreeAnalyzer::Analyze()76+void TreeAnalyzer::Analyze() { DeepFirstSearch(root_); }
71-{
72- DeepFirstSearch(root_);
73-}
74 77 
75-HostTasks &TreeAnalyzer::GetHCCLTasks()78+HostTasks &TreeAnalyzer::GetHCCLTasks() { return hcclTasks_; }
76-{
77- return hcclTasks_;
78-}
79 79 
80-HostTasks &TreeAnalyzer::GetComputeTasks()80+HostTasks &TreeAnalyzer::GetComputeTasks() { return computeTasks_; }
81-{
82- return computeTasks_;
83-}
84 81 
85-HostTasks &TreeAnalyzer::GetTasks()82+HostTasks &TreeAnalyzer::GetTasks() { return tasks_; }
86-{
87- return tasks_;
88-}
89 83 
90-HCCLBigOpDescs &TreeAnalyzer::GetHcclBigOps()84+HCCLBigOpDescs &TreeAnalyzer::GetHcclBigOps() { return hcclBigOpDescs_; }
91-{
92- return hcclBigOpDescs_;
93-}
94 85 
95-GeFusionOpInfos &TreeAnalyzer::GetGeFusionOpInfos()86+GeFusionOpInfos &TreeAnalyzer::GetGeFusionOpInfos() { return geFusionOpInfos_; }
96-{
97- return geFusionOpInfos_;
98-}
99 87 
100void TreeAnalyzer::DeepFirstSearch(const std::shared_ptr<TreeNode> &node)88void TreeAnalyzer::DeepFirstSearch(const std::shared_ptr<TreeNode> &node)
101{89{
102- if (!node || !node->event) {90+ if (!node || !node->event)
91+ {
103 ERROR("Node pointer or event is nullptr, threadId = %", threadId_);92 ERROR("Node pointer or event is nullptr, threadId = %", threadId_);
104 return;93 return;
105 }94 }
106 AnalyzeNode(node);95 AnalyzeNode(node);
107 path_[node->event->info.level] = node;96 path_[node->event->info.level] = node;
108 // DFS97 // DFS
109- for (const auto &child: node->children) {98+ for (const auto &child : node->children)
99+ {
110 DeepFirstSearch(child);100 DeepFirstSearch(child);
111- if (node->event->info.level == child->event->info.level) {101+ if (node->event->info.level == child->event->info.level)
102+ {
112 path_[node->event->info.level] = node;103 path_[node->event->info.level] = node;
113 }104 }
114 }105 }
@@ -117,9 +108,12 @@ void TreeAnalyzer::DeepFirstSearch(const std::shared_ptr<TreeNode> &node)
117 108 
118void TreeAnalyzer::AnalyzeNode(const std::shared_ptr<TreeNode> &node)109void TreeAnalyzer::AnalyzeNode(const std::shared_ptr<TreeNode> &node)
119{110{
120- if (node->event->info.level == MSPROF_REPORT_RUNTIME_LEVEL) {111+ if (node->event->info.level == MSPROF_REPORT_RUNTIME_LEVEL)
112+ {
121 AnalyzeRuntimeNode(node);113 AnalyzeRuntimeNode(node);
122- } else if (node->event->info.level == MSPROF_REPORT_MODEL_LEVEL) {114+ }
115+ else if (node->event->info.level == MSPROF_REPORT_MODEL_LEVEL)
116+ {
123 AnalyzeModelNode(node);117 AnalyzeModelNode(node);
124 }118 }
125}119}
@@ -128,10 +122,13 @@ void TreeAnalyzer::AnalyzeRuntimeNode(const std::shared_ptr<TreeNode> &node)
128{122{
129 auto isHccl = IsHcclTask();123 auto isHccl = IsHcclTask();
130 auto isOperateCompute = IsComputeTask(node);124 auto isOperateCompute = IsComputeTask(node);
131- if (isOperateCompute) {125+ if (isOperateCompute)
126+ {
132 auto computeTasks = GetComputeTaskDescs(node);127 auto computeTasks = GetComputeTaskDescs(node);
133- for (auto &task : computeTasks) {128+ for (auto &task : computeTasks)
134- if (task->op && task->op->type == OpType::OPTYPE_INVALID) {129+ {
130+ if (task->op && task->op->type == OpType::OPTYPE_INVALID)
131+ {
135 // 该任务不会被认为是compute任务132 // 该任务不会被认为是compute任务
136 continue;133 continue;
137 }134 }
@@ -141,19 +138,24 @@ void TreeAnalyzer::AnalyzeRuntimeNode(const std::shared_ptr<TreeNode> &node)
141 tasks_.insert(tasks_.end(), computeTasks.begin(), computeTasks.end());138 tasks_.insert(tasks_.end(), computeTasks.begin(), computeTasks.end());
142 }139 }
143 140 
144- if (isHccl) {141+ if (isHccl)
142+ {
145 auto hcclTasks = GetHcclTaskDescs(node);143 auto hcclTasks = GetHcclTaskDescs(node);
146 144 
147- for (auto &task : hcclTasks) {145+ for (auto &task : hcclTasks)
148- if (task->op && task->op->type == OpType::OPTYPE_INVALID) {146+ {
147+ if (task->op && task->op->type == OpType::OPTYPE_INVALID)
148+ {
149 // 该任务不会被认为是hccl任务149 // 该任务不会被认为是hccl任务
150 continue;150 continue;
151 }151 }
152 hcclTasks_.emplace_back(task);152 hcclTasks_.emplace_back(task);
153 }153 }
154 154 
155- for (auto &task : hcclTasks) {155+ for (auto &task : hcclTasks)
156- if (isOperateCompute) {156+ {
157+ if (isOperateCompute)
158+ {
157 // task已插入，无需再次插入, ffts+模式的通信算子不会进入该分支159 // task已插入，无需再次插入, ffts+模式的通信算子不会进入该分支
158 continue;160 continue;
159 }161 }
@@ -162,13 +164,16 @@ void TreeAnalyzer::AnalyzeRuntimeNode(const std::shared_ptr<TreeNode> &node)
162 UpdateHcclBigOpDescs(node);164 UpdateHcclBigOpDescs(node);
163 }165 }
164 166 
165- if (!isOperateCompute && !isHccl) {167+ if (!isOperateCompute && !isHccl)
168+ {
166 auto otherTask = GetOtherTaskDesc(node);169 auto otherTask = GetOtherTaskDesc(node);
167- if (otherTask) {170+ if (otherTask)
171+ {
168 tasks_.emplace_back(otherTask);172 tasks_.emplace_back(otherTask);
169 auto taskType = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, otherTask->taskType);173 auto taskType = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, otherTask->taskType);
170 // 对于纯rts_track数据,只有算子类型在白名单中才生成computeTask数据174 // 对于纯rts_track数据,只有算子类型在白名单中才生成computeTask数据
171- if (KERNEL_COMPUTE_WHITE_LIST.find(taskType) != KERNEL_COMPUTE_WHITE_LIST.end()) {175+ if (KERNEL_COMPUTE_WHITE_LIST.find(taskType) != KERNEL_COMPUTE_WHITE_LIST.end())
176+ {
172 computeTasks_.emplace_back(otherTask);177 computeTasks_.emplace_back(otherTask);
173 }178 }
174 }179 }
@@ -179,14 +184,18 @@ void TreeAnalyzer::AnalyzeModelNode(const std::shared_ptr<TreeNode> &node)
179{184{
180 auto modelApi = (node != nullptr && node->event != nullptr) ? node->event->apiPtr : nullptr;185 auto modelApi = (node != nullptr && node->event != nullptr) ? node->event->apiPtr : nullptr;
181 std::string key = Utils::Join("_", node->event->info.start, node->event->info.end);186 std::string key = Utils::Join("_", node->event->info.start, node->event->info.end);
182- if (visitedModel_.find(key) != visitedModel_.end()) {187+ if (visitedModel_.find(key) != visitedModel_.end())
188+ {
183 return;189 return;
184 }190 }
185 auto modelId = modelApi != nullptr ? modelApi->itemId : INVALID_MODEL_ID;191 auto modelId = modelApi != nullptr ? modelApi->itemId : INVALID_MODEL_ID;
186- for (const auto &record: node->records) {192+ for (const auto &record : node->records)
187- if (record->info.type == EventType::EVENT_TYPE_FUSION_OP_INFO) {193+ {
194+ if (record->info.type == EventType::EVENT_TYPE_FUSION_OP_INFO)
195+ {
188 auto fusionOp = record->additionPtr;196 auto fusionOp = record->additionPtr;
189- if (fusionOp) {197+ if (fusionOp)
198+ {
190 auto fusionStruct = Utils::ReinterpretConvert<ProfFusionOpInfo *>(fusionOp->data);199 auto fusionStruct = Utils::ReinterpretConvert<ProfFusionOpInfo *>(fusionOp->data);
191 std::shared_ptr<ProfFusionOpInfo> fusionOpInfo;200 std::shared_ptr<ProfFusionOpInfo> fusionOpInfo;
192 MAKE_SHARED_RETURN_VOID(fusionOpInfo, ProfFusionOpInfo, *fusionStruct);201 MAKE_SHARED_RETURN_VOID(fusionOpInfo, ProfFusionOpInfo, *fusionStruct);
@@ -199,40 +208,44 @@ void TreeAnalyzer::AnalyzeModelNode(const std::shared_ptr<TreeNode> &node)
199 visitedModel_.insert(key);208 visitedModel_.insert(key);
200}209}
201 210 
202-bool TreeAnalyzer::IsHcclTask()211+bool TreeAnalyzer::IsHcclTask() { return path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) != path_.end(); }
203-{
204- return path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) != path_.end();
205-}
206 212 
207bool TreeAnalyzer::IsComputeTask(const std::shared_ptr<TreeNode> &node)213bool TreeAnalyzer::IsComputeTask(const std::shared_ptr<TreeNode> &node)
208{214{
209- if (path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end()) {215+ if (path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end())
216+ {
210 return false;217 return false;
211 }218 }
212 219 
213 // 临时规避 lccl算子过滤,后续依据lccl 特征独立判定220 // 临时规避 lccl算子过滤,后续依据lccl 特征独立判定
214 auto node_api = path_[MSPROF_REPORT_NODE_LEVEL]->event->apiPtr;221 auto node_api = path_[MSPROF_REPORT_NODE_LEVEL]->event->apiPtr;
215- if (!node_api || TypeData::GetInstance().Get(node_api->level, node_api->type) == GE_STEP_INFO_API_TYPE) {222+ if (!node_api || TypeData::GetInstance().Get(node_api->level, node_api->type) == GE_STEP_INFO_API_TYPE)
223+ {
216 return false;224 return false;
217 }225 }
218 std::string itemId = HashData::GetInstance().Get(node_api->itemId);226 std::string itemId = HashData::GetInstance().Get(node_api->itemId);
219- if (itemId.substr(0, LCCL_PREFIX.size()) == LCCL_PREFIX) {227+ if (itemId.substr(0, LCCL_PREFIX.size()) == LCCL_PREFIX)
228+ {
220 return false;229 return false;
221 }230 }
222 231 
223- for (const auto &record: node->records) {232+ for (const auto &record : node->records)
224- if (record->info.type != EventType::EVENT_TYPE_TASK_TRACK) {233+ {
234+ if (record->info.type != EventType::EVENT_TYPE_TASK_TRACK)
235+ {
225 continue;236 continue;
226 }237 }
227 238 
228 auto trace = record->compactPtr;239 auto trace = record->compactPtr;
229- auto taskType = TypeData::GetInstance().Get(node->event->info.level,240+ auto taskType = TypeData::GetInstance().Get(node->event->info.level, trace->data.runtimeTrack.taskType);
230- trace->data.runtimeTrack.taskType);
231 if (taskType.substr(0, KERNEL_TASK_PREFIX.size()) == KERNEL_TASK_PREFIX ||241 if (taskType.substr(0, KERNEL_TASK_PREFIX.size()) == KERNEL_TASK_PREFIX ||
232- taskType == KERNEL_STARS_COMMON_TASK_TYPE) {242+ taskType == KERNEL_STARS_COMMON_TASK_TYPE)
243+ {
233 // 传统core task和dsa task244 // 传统core task和dsa task
234 return true;245 return true;
235- } else if (taskType == KERNEL_FFTS_PLUS_TASK_TYPE) {246+ }
247+ else if (taskType == KERNEL_FFTS_PLUS_TASK_TYPE)
248+ {
236 // ffts+ task249 // ffts+ task
237 return path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) == path_.end();250 return path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) == path_.end();
238 }251 }
@@ -243,55 +256,60 @@ bool TreeAnalyzer::IsComputeTask(const std::shared_ptr<TreeNode> &node)
243 256 
244void TreeAnalyzer::UpdateHcclBigOpDescs(const std::shared_ptr<TreeNode> &node)257void TreeAnalyzer::UpdateHcclBigOpDescs(const std::shared_ptr<TreeNode> &node)
245{258{
246- if (path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end()) {259+ if (path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end())
260+ {
247 return;261 return;
248 }262 }
249 auto nodeNode = path_[MSPROF_REPORT_NODE_LEVEL];263 auto nodeNode = path_[MSPROF_REPORT_NODE_LEVEL];
250 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);264 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);
251- if (tracks.empty()) {265+ if (tracks.empty())
266+ {
252 ERROR("TreeNode task track records is empty, threadId = %", threadId_);267 ERROR("TreeNode task track records is empty, threadId = %", threadId_);
253 return;268 return;
254 }269 }
255 auto track = tracks.back()->compactPtr;270 auto track = tracks.back()->compactPtr;
256 auto deviceId = track->data.runtimeTrack.deviceId;271 auto deviceId = track->data.runtimeTrack.deviceId;
257 272 
258- if (!hcclBigOpDescs_.empty() &&273+ if (!hcclBigOpDescs_.empty() && hcclBigOpDescs_.back()->hcclBigOpDesc->endTime == nodeNode->event->info.end)
259- hcclBigOpDescs_.back()->hcclBigOpDesc->endTime == nodeNode->event->info.end) {274+ {
260 WARN("Find same end time when update Hccl big op descs, threadId = %", threadId_);275 WARN("Find same end time when update Hccl big op descs, threadId = %", threadId_);
261 return;276 return;
262 }277 }
263 278 
264- auto modelTrace = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ?279+ auto modelTrace = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ? path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
265- path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
266 auto modelApi = modelTrace != nullptr ? modelTrace->event->apiPtr : nullptr;280 auto modelApi = modelTrace != nullptr ? modelTrace->event->apiPtr : nullptr;
267 auto indexId = modelApi != nullptr ? modelApi->reserve : -1;281 auto indexId = modelApi != nullptr ? modelApi->reserve : -1;
268 auto modelId = GetModelId(modelApi, deviceId, track->data.runtimeTrack.streamId,282 auto modelId = GetModelId(modelApi, deviceId, track->data.runtimeTrack.streamId,
269- static_cast<uint16_t>(track->data.runtimeTrack.taskId >> TWO_BYTES),283+ static_cast<uint16_t>(track->data.runtimeTrack.taskId >> TWO_BYTES), track->timeStamp);
270- track->timeStamp);
271 auto connectionId = nodeNode->event->id;284 auto connectionId = nodeNode->event->id;
272 auto nodeRecords = GetNodeRecordsByType(nodeNode, EventType::EVENT_TYPE_NODE_BASIC_INFO);285 auto nodeRecords = GetNodeRecordsByType(nodeNode, EventType::EVENT_TYPE_NODE_BASIC_INFO);
273 std::shared_ptr<MsprofCompactInfo> nodeDesc = nullptr;286 std::shared_ptr<MsprofCompactInfo> nodeDesc = nullptr;
274- if (!nodeRecords.empty() && nodeRecords.front() != nullptr) {287+ if (!nodeRecords.empty() && nodeRecords.front() != nullptr)
288+ {
275 nodeDesc = nodeRecords.front()->compactPtr;289 nodeDesc = nodeRecords.front()->compactPtr;
276 }290 }
277 auto hcclOpRecords = GetNodeRecordsByType(nodeNode, EventType::EVENT_TYPE_HCCL_OP_INFO);291 auto hcclOpRecords = GetNodeRecordsByType(nodeNode, EventType::EVENT_TYPE_HCCL_OP_INFO);
278 std::shared_ptr<MsprofCompactInfo> hcclOpDesc = nullptr;292 std::shared_ptr<MsprofCompactInfo> hcclOpDesc = nullptr;
279- if (!hcclOpRecords.empty() && hcclOpRecords.front() != nullptr) {293+ if (!hcclOpRecords.empty() && hcclOpRecords.front() != nullptr)
294+ {
280 hcclOpDesc = hcclOpRecords.front()->compactPtr;295 hcclOpDesc = hcclOpRecords.front()->compactPtr;
281- } else {296+ }
297+ else
298+ {
282 ERROR("Not report hccl op info for api: %", modelId);299 ERROR("Not report hccl op info for api: %", modelId);
283 }300 }
284 int64_t kfcConnectionId = INVALID_VALUE;301 int64_t kfcConnectionId = INVALID_VALUE;
285- for (const auto &child: nodeNode->children) {302+ for (const auto &child : nodeNode->children)
286- if (child->event->info.level == MSPROF_REPORT_NODE_LEVEL) {303+ {
304+ if (child->event->info.level == MSPROF_REPORT_NODE_LEVEL)
305+ {
287 kfcConnectionId = child->event->id;306 kfcConnectionId = child->event->id;
288 }307 }
289 }308 }
290 309 
291 auto nodeApi = nodeNode->event->apiPtr;310 auto nodeApi = nodeNode->event->apiPtr;
292 std::shared_ptr<HcclBigOpDesc> desc;311 std::shared_ptr<HcclBigOpDesc> desc;
293- MAKE_SHARED_RETURN_VOID(desc, HcclBigOpDesc, nodeApi->beginTime,312+ MAKE_SHARED_RETURN_VOID(desc, HcclBigOpDesc, nodeApi->beginTime, nodeApi->endTime, deviceId, modelId, indexId,
294- nodeApi->endTime, deviceId, modelId, indexId,
295 connectionId, track->threadId, nodeDesc, hcclOpDesc, kfcConnectionId);313 connectionId, track->threadId, nodeDesc, hcclOpDesc, kfcConnectionId);
296 std::shared_ptr<Operator> op;314 std::shared_ptr<Operator> op;
297 MAKE_SHARED_RETURN_VOID(op, Operator, desc, nodeApi->itemId, OpType::OPTYPE_HCCL_BIG);315 MAKE_SHARED_RETURN_VOID(op, Operator, desc, nodeApi->itemId, OpType::OPTYPE_HCCL_BIG);
@@ -302,8 +320,10 @@ std::vector<std::shared_ptr<Event>> TreeAnalyzer::GetNodeRecordsByType(const std
302 const EventType &type)320 const EventType &type)
303{321{
304 std::vector<std::shared_ptr<Event>> records;322 std::vector<std::shared_ptr<Event>> records;
305- for (const auto &record: node->records) {323+ for (const auto &record : node->records)
306- if (record->info.type == type) {324+ {
325+ if (record->info.type == type)
326+ {
307 records.emplace_back(record);327 records.emplace_back(record);
308 }328 }
309 }329 }
@@ -312,9 +332,8 @@ std::vector<std::shared_ptr<Event>> TreeAnalyzer::GetNodeRecordsByType(const std
312 332 
313std::shared_ptr<HostTask> TreeAnalyzer::GenHostTask(const std::shared_ptr<MsprofCompactInfo> &track,333std::shared_ptr<HostTask> TreeAnalyzer::GenHostTask(const std::shared_ptr<MsprofCompactInfo> &track,
314 const std::shared_ptr<MsprofApi> &modelApi,334 const std::shared_ptr<MsprofApi> &modelApi,
315- const std::shared_ptr<Operator> &opPtr,335+ const std::shared_ptr<Operator> &opPtr, uint32_t ctxId,
316- uint32_t ctxId, uint16_t taskType,336+ uint16_t taskType, int64_t connectionId)
317- int64_t connectionId)
318{337{
319 std::shared_ptr<HostTask> task;338 std::shared_ptr<HostTask> task;
320 MAKE_SHARED0_RETURN_VALUE(task, HostTask, nullptr);339 MAKE_SHARED0_RETURN_VALUE(task, HostTask, nullptr);
@@ -335,51 +354,59 @@ std::shared_ptr<HostTask> TreeAnalyzer::GenHostTask(const std::shared_ptr<Msprof
335 return task;354 return task;
336}355}
337 356 
338-HostTasks TreeAnalyzer::GenComputeHostTasks(ComputeOpDescs &ops,357+HostTasks TreeAnalyzer::GenComputeHostTasks(ComputeOpDescs &ops, const std::shared_ptr<MsprofCompactInfo> &track,
339- const std::shared_ptr<MsprofCompactInfo> &track,
340 int64_t connection_id)358 int64_t connection_id)
341{359{
342- auto modelNode = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ?360+ auto modelNode = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ? path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
343- path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
344 auto modelApi = modelNode == nullptr ? nullptr : modelNode->event->apiPtr;361 auto modelApi = modelNode == nullptr ? nullptr : modelNode->event->apiPtr;
345 // L0场景没有上报Node层补充信息且该任务非ctx类任务362 // L0场景没有上报Node层补充信息且该任务非ctx类任务
346- if (ops.empty()) {363+ if (ops.empty())
364+ {
347 // 补充一条有效的Op信息365 // 补充一条有效的Op信息
348- auto nodeNode = path_.find(MSPROF_REPORT_NODE_LEVEL) != path_.end() ?366+ auto nodeNode = path_.find(MSPROF_REPORT_NODE_LEVEL) != path_.end() ? path_[MSPROF_REPORT_NODE_LEVEL] : nullptr;
349- path_[MSPROF_REPORT_NODE_LEVEL] : nullptr;367+ if (nodeNode == nullptr)
350- if (nodeNode == nullptr) {368+ {
351 return HostTasks{};369 return HostTasks{};
352 }370 }
353 std::shared_ptr<Operator> op;371 std::shared_ptr<Operator> op;
354 std::shared_ptr<OpDesc> desc;372 std::shared_ptr<OpDesc> desc;
355- MAKE_SHARED_RETURN_VALUE(op, Operator, {}, desc, nodeNode->event->apiPtr->itemId, OpType::OPTYPE_RESERVED);373+ MAKE_SHARED_RETURN_VALUE(desc, OpDesc, {});
356- auto task = GenHostTask(track, modelApi, op,374+ desc->runtimeTrackDesc = track;
357- DEFAULT_CONTEXT_ID, track->data.runtimeTrack.taskType, connection_id);375+ MAKE_SHARED_RETURN_VALUE(op, Operator, {}, desc, nodeNode->event->apiPtr->itemId, OpType::OPTYPE_COMPUTE);
376+ auto task =
377+ GenHostTask(track, modelApi, op, DEFAULT_CONTEXT_ID, track->data.runtimeTrack.taskType, connection_id);
358 return (task != nullptr) ? HostTasks{task} : HostTasks{};378 return (task != nullptr) ? HostTasks{task} : HostTasks{};
359 }379 }
360 380 
361 HostTasks results;381 HostTasks results;
362- for (const auto &pair: ops) {382+ for (const auto &pair : ops)
383+ {
363 auto desc = pair.second->opDesc;384 auto desc = pair.second->opDesc;
364- if (!desc) {385+ if (!desc)
365- ERROR("No op % desc found for task timestamp: % when gen compute task",386+ {
366- pair.second->name, track->timeStamp);387+ ERROR("No op % desc found for task timestamp: % when gen compute task", pair.second->name,
388+ track->timeStamp);
367 continue;389 continue;
368 }390 }
369 391 
370 std::vector<uint32_t> ctxIds;392 std::vector<uint32_t> ctxIds;
371 // 只有FFTS+类型应该保留CtxId393 // 只有FFTS+类型应该保留CtxId
372- if (desc->ctxId) {394+ if (desc->ctxId)
395+ {
373 auto ctxIdInfo = ReinterpretConvert<MsprofContextIdInfo *>(desc->ctxId->data);396 auto ctxIdInfo = ReinterpretConvert<MsprofContextIdInfo *>(desc->ctxId->data);
374 ctxIds.assign(ctxIdInfo->ctxIds, ctxIdInfo->ctxIds + ctxIdInfo->ctxIdNum);397 ctxIds.assign(ctxIdInfo->ctxIds, ctxIdInfo->ctxIds + ctxIdInfo->ctxIdNum);
375- } else {398+ }
399+ else
400+ {
376 ctxIds = {DEFAULT_CONTEXT_ID};401 ctxIds = {DEFAULT_CONTEXT_ID};
377 }402 }
378- 403+ desc->runtimeTrackDesc = track;
379- for (const auto &ctxId: ctxIds) {404+ for (const auto &ctxId : ctxIds)
380- auto task = GenHostTask(track, modelApi, pair.second, ctxId,405+ {
381- track->data.runtimeTrack.taskType, connection_id);406+ auto task =
382- if (task) {407+ GenHostTask(track, modelApi, pair.second, ctxId, track->data.runtimeTrack.taskType, connection_id);
408+ if (task)
409+ {
383 results.emplace_back(task);410 results.emplace_back(task);
384 }411 }
385 }412 }
@@ -387,36 +414,40 @@ HostTasks TreeAnalyzer::GenComputeHostTasks(ComputeOpDescs &ops,
387 return results;414 return results;
388}415}
389 416 
390-void TreeAnalyzer::UpdateComputeDescForFftsSituation(ComputeOpDescs &descs,417+void TreeAnalyzer::UpdateComputeDescForFftsSituation(ComputeOpDescs &descs, const std::shared_ptr<Event> &track)
391- const std::shared_ptr<Event> &track)
392{418{
393 std::string specialKey;419 std::string specialKey;
394 // notice： 特殊场景，刷新op描述420 // notice： 特殊场景，刷新op描述
395- for (auto &pair : descs) {421+ for (auto &pair : descs)
396- if (!pair.second->opDesc) {422+ {
423+ if (!pair.second->opDesc)
424+ {
397 ERROR("Illegal compute op, no op desc found, threadId = %", threadId_);425 ERROR("Illegal compute op, no op desc found, threadId = %", threadId_);
398 continue;426 continue;
399 }427 }
400- if (!pair.second->opDesc->nodeDesc) {428+ if (!pair.second->opDesc->nodeDesc)
429+ {
401 continue;430 continue;
402 }431 }
403 // 特殊场景1： ffts+模式会上报一条多余的nodebasic用于标记ffts图，删除432 // 特殊场景1： ffts+模式会上报一条多余的nodebasic用于标记ffts图，删除
404- if (pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType == FFTS_PLUS_TASK_TYPE) {433+ if (pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType == FFTS_PLUS_TASK_TYPE)
434+ {
405 specialKey = pair.first;435 specialKey = pair.first;
406 // 该数据在一个task中只有1条436 // 该数据在一个task中只有1条
407 break;437 break;
408 }438 }
409 }439 }
410- if (!specialKey.empty()) {440+ if (!specialKey.empty())
441+ {
411 descs.erase(specialKey);442 descs.erase(specialKey);
412 }443 }
413}444}
414 445 
415-void TreeAnalyzer::UpdateComputeDescForHcclSituation(ComputeOpDescs &descs,446+void TreeAnalyzer::UpdateComputeDescForHcclSituation(ComputeOpDescs &descs, const std::shared_ptr<Event> &track,
416- const std::shared_ptr<Event> &track,
417 uint64_t item_id)447 uint64_t item_id)
418{448{
419- if (descs.empty()) {449+ if (descs.empty())
450+ {
420 // 补充一个临时的算子描述，只发生在L0场景451 // 补充一个临时的算子描述，只发生在L0场景
421 std::shared_ptr<OpDesc> desc;452 std::shared_ptr<OpDesc> desc;
422 std::shared_ptr<Operator> op;453 std::shared_ptr<Operator> op;
@@ -425,12 +456,15 @@ void TreeAnalyzer::UpdateComputeDescForHcclSituation(ComputeOpDescs &descs,
425 auto name = Utils::Join("_", std::to_string(PLACEHOLDER_OP_NAME), std::to_string(UINT64_MAX));456 auto name = Utils::Join("_", std::to_string(PLACEHOLDER_OP_NAME), std::to_string(UINT64_MAX));
426 descs[name] = op;457 descs[name] = op;
427 }458 }
428- for (auto &pair : descs) {459+ for (auto &pair : descs)
429- if (!pair.second->opDesc) {460+ {
461+ if (!pair.second->opDesc)
462+ {
430 ERROR("Illegal compute op, no op desc found, threadId = %", threadId_);463 ERROR("Illegal compute op, no op desc found, threadId = %", threadId_);
431 continue;464 continue;
432 }465 }
433- if (!pair.second->opDesc->nodeDesc) {466+ if (!pair.second->opDesc->nodeDesc)
467+ {
434 // 避免内存泄露468 // 避免内存泄露
435 MAKE_SHARED0_NO_OPERATION(pair.second->opDesc->nodeDesc, MsprofCompactInfo);469 MAKE_SHARED0_NO_OPERATION(pair.second->opDesc->nodeDesc, MsprofCompactInfo);
436 }470 }
@@ -438,10 +472,13 @@ void TreeAnalyzer::UpdateComputeDescForHcclSituation(ComputeOpDescs &descs,
438 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.opName = item_id;472 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.opName = item_id;
439 auto taskType = track->compactPtr->data.runtimeTrack.taskType;473 auto taskType = track->compactPtr->data.runtimeTrack.taskType;
440 auto taskTypeStr = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, taskType);474 auto taskTypeStr = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, taskType);
441- if (taskTypeStr == KERNEL_AI_CPU_TASK_TYPE) {475+ if (taskTypeStr == KERNEL_AI_CPU_TASK_TYPE)
476+ {
442 // 特殊场景3：HCCL AICPU任务，该任务需要将任务类型刷为HCCL_AICPU477 // 特殊场景3：HCCL AICPU任务，该任务需要将任务类型刷为HCCL_AICPU
443 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType = HCCL_AI_CPU_TASK_TYPE;478 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType = HCCL_AI_CPU_TASK_TYPE;
444- } else if (taskTypeStr == KERNEL_AI_CORE_TASK_TYPE) {479+ }
480+ else if (taskTypeStr == KERNEL_AI_CORE_TASK_TYPE)
481+ {
445 // 特殊场景4：HCCL AICORE(Reduce TBE)任务，该任务需要将任务类型刷为AI_CORE482 // 特殊场景4：HCCL AICORE(Reduce TBE)任务，该任务需要将任务类型刷为AI_CORE
446 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType = HCCL_TASK_TYPE;483 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType = HCCL_TASK_TYPE;
447 }484 }
@@ -450,12 +487,15 @@ void TreeAnalyzer::UpdateComputeDescForHcclSituation(ComputeOpDescs &descs,
450 487 
451void TreeAnalyzer::UpdateComputeDescForHelperSituation(ComputeOpDescs &descs)488void TreeAnalyzer::UpdateComputeDescForHelperSituation(ComputeOpDescs &descs)
452{489{
453- for (auto &pair : descs) {490+ for (auto &pair : descs)
454- if (!pair.second->opDesc || !pair.second->opDesc->nodeDesc) {491+ {
492+ if (!pair.second->opDesc || !pair.second->opDesc->nodeDesc)
493+ {
455 continue;494 continue;
456 }495 }
457 // helper场景：HCCL算子运行在AI_CPU上, 但没有HCCL层api，任务类型刷为AICPU496 // helper场景：HCCL算子运行在AI_CPU上, 但没有HCCL层api，任务类型刷为AICPU
458- if (pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType == HCCL_TASK_TYPE) {497+ if (pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType == HCCL_TASK_TYPE)
498+ {
459 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType = AICPU_TASK_TYPE;499 pair.second->opDesc->nodeDesc->data.nodeBasicInfo.taskType = AICPU_TASK_TYPE;
460 }500 }
461 }501 }
@@ -463,32 +503,35 @@ void TreeAnalyzer::UpdateComputeDescForHelperSituation(ComputeOpDescs &descs)
463 503 
464HostTasks TreeAnalyzer::GetComputeTaskDescs(const std::shared_ptr<TreeNode> &node)504HostTasks TreeAnalyzer::GetComputeTaskDescs(const std::shared_ptr<TreeNode> &node)
465{505{
466- if (path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end()) {506+ if (path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end())
507+ {
467 return {};508 return {};
468 }509 }
469- auto nodeNode = path_.find(MSPROF_REPORT_NODE_LEVEL) != path_.end() ?510+ auto nodeNode = path_.find(MSPROF_REPORT_NODE_LEVEL) != path_.end() ? path_[MSPROF_REPORT_NODE_LEVEL] : nullptr;
470- path_[MSPROF_REPORT_NODE_LEVEL] : nullptr;511+ auto modelNode = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ? path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
471- auto modelNode = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ?
472- path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
473 512 
474 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);513 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);
475- if (tracks.empty()) {514+ if (tracks.empty())
515+ {
476 ERROR("TreeNode task track records is empty, threadId = %", threadId_);516 ERROR("TreeNode task track records is empty, threadId = %", threadId_);
477 return {};517 return {};
478 }518 }
479 519 
480- std::string type = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL,520+ std::string type =
481- tracks.back()->compactPtr->data.runtimeTrack.taskType);521+ TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, tracks.back()->compactPtr->data.runtimeTrack.taskType);
482 bool useCtxId = CONTEXT_ID_WHITE_LIST.find(type) != CONTEXT_ID_WHITE_LIST.end();522 bool useCtxId = CONTEXT_ID_WHITE_LIST.find(type) != CONTEXT_ID_WHITE_LIST.end();
483 ComputeOpDescs ops = GetComputeOpDescs(nodeNode, useCtxId);523 ComputeOpDescs ops = GetComputeOpDescs(nodeNode, useCtxId);
484 // 特殊场景，刷新算子描述524 // 特殊场景，刷新算子描述
485 UpdateComputeDescForFftsSituation(ops, tracks.back());525 UpdateComputeDescForFftsSituation(ops, tracks.back());
486- if (IsHcclTask()) {526+ if (IsHcclTask())
487- auto hcclNode = path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) != path_.end() ?527+ {
488- path_[MSPROF_REPORT_HCCL_NODE_LEVEL] : nullptr;528+ auto hcclNode =
529+ path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) != path_.end() ? path_[MSPROF_REPORT_HCCL_NODE_LEVEL] : nullptr;
489 auto item_id = hcclNode == nullptr ? 0 : hcclNode->event->apiPtr->itemId;530 auto item_id = hcclNode == nullptr ? 0 : hcclNode->event->apiPtr->itemId;
490 UpdateComputeDescForHcclSituation(ops, tracks.back(), item_id);531 UpdateComputeDescForHcclSituation(ops, tracks.back(), item_id);
491- } else if (type == KERNEL_AI_CPU_TASK_TYPE) {532+ }
533+ else if (type == KERNEL_AI_CPU_TASK_TYPE)
534+ {
492 UpdateComputeDescForHelperSituation(ops);535 UpdateComputeDescForHelperSituation(ops);
493 }536 }
494 auto track = tracks.back()->compactPtr;537 auto track = tracks.back()->compactPtr;
@@ -499,19 +542,18 @@ HostTasks TreeAnalyzer::GetComputeTaskDescs(const std::shared_ptr<TreeNode> &nod
499 542 
500HostTasks TreeAnalyzer::GetHcclTaskDescs(const std::shared_ptr<TreeNode> &node)543HostTasks TreeAnalyzer::GetHcclTaskDescs(const std::shared_ptr<TreeNode> &node)
501{544{
502- if (path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) == path_.end() ||545+ if (path_.find(MSPROF_REPORT_HCCL_NODE_LEVEL) == path_.end() || path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end())
503- path_.find(MSPROF_REPORT_NODE_LEVEL) == path_.end()) {546+ {
504 return {};547 return {};
505 }548 }
506 auto hcclNode = path_[MSPROF_REPORT_HCCL_NODE_LEVEL];549 auto hcclNode = path_[MSPROF_REPORT_HCCL_NODE_LEVEL];
507 auto nodeNode = path_[MSPROF_REPORT_NODE_LEVEL];550 auto nodeNode = path_[MSPROF_REPORT_NODE_LEVEL];
508- auto modelTrace = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ?551+ auto modelTrace = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ? path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
509- path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;552+ auto modelApi = modelTrace != nullptr ? modelTrace->event->apiPtr : nullptr;
510- auto modelApi = modelTrace != nullptr ?
511- modelTrace->event->apiPtr : nullptr;
512 553 
513 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);554 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);
514- if (tracks.empty()) {555+ if (tracks.empty())
556+ {
515 ERROR("TreeNode task track records is empty, threadId = %", threadId_);557 ERROR("TreeNode task track records is empty, threadId = %", threadId_);
516 return {};558 return {};
517 }559 }
@@ -520,16 +562,19 @@ HostTasks TreeAnalyzer::GetHcclTaskDescs(const std::shared_ptr<TreeNode> &node)
520 562 
521 HostTasks results;563 HostTasks results;
522 auto ret = Utils::Reserve(results, hcclOpDescs.size());564 auto ret = Utils::Reserve(results, hcclOpDescs.size());
523- if (!ret) {565+ if (!ret)
566+ {
524 ERROR("Reserve results failed, threadId = %", threadId_);567 ERROR("Reserve results failed, threadId = %", threadId_);
525 return results;568 return results;
526 }569 }
527 570 
528- for (const auto &pair: hcclOpDescs) {571+ for (const auto &pair : hcclOpDescs)
572+ {
529 auto desc = pair.second->hcclSmallOpDesc;573 auto desc = pair.second->hcclSmallOpDesc;
530- auto task = GenHostTask(track, modelApi, pair.second, desc->ctxId,574+ auto task = GenHostTask(track, modelApi, pair.second, desc->ctxId, track->data.runtimeTrack.taskType,
531- track->data.runtimeTrack.taskType, nodeNode->event->id);575+ nodeNode->event->id);
532- if (task) {576+ if (task)
577+ {
533 results.emplace_back(task);578 results.emplace_back(task);
534 }579 }
535 }580 }
@@ -538,69 +583,82 @@ HostTasks TreeAnalyzer::GetHcclTaskDescs(const std::shared_ptr<TreeNode> &node)
538 583 
539std::shared_ptr<HostTask> TreeAnalyzer::GetOtherTaskDesc(const std::shared_ptr<TreeNode> &node)584std::shared_ptr<HostTask> TreeAnalyzer::GetOtherTaskDesc(const std::shared_ptr<TreeNode> &node)
540{585{
541- auto modelNode = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ?586+ auto modelNode = path_.find(MSPROF_REPORT_MODEL_LEVEL) != path_.end() ? path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;
542- path_[MSPROF_REPORT_MODEL_LEVEL] : nullptr;587+ auto modelApi = modelNode != nullptr ? modelNode->event->apiPtr : nullptr;
543- auto modelApi = modelNode != nullptr ?
544- modelNode->event->apiPtr : nullptr;
545 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);588 auto tracks = GetNodeRecordsByType(node, EventType::EVENT_TYPE_TASK_TRACK);
546- if (tracks.empty()) {589+ if (tracks.empty())
590+ {
547 ERROR("TreeNode task track records is empty, threadId = %", threadId_);591 ERROR("TreeNode task track records is empty, threadId = %", threadId_);
548 return nullptr;592 return nullptr;
549 }593 }
550 auto track = tracks.back()->compactPtr;594 auto track = tracks.back()->compactPtr;
551 auto taskType = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, track->data.runtimeTrack.taskType);595 auto taskType = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, track->data.runtimeTrack.taskType);
552 uint32_t contextId =596 uint32_t contextId =
553- (taskType == KERNEL_MIX_AIC_TASK_TYPE || taskType == KERNEL_MIX_AIV_TASK_TYPE) ? 0 : DEFAULT_CONTEXT_ID;597+ (taskType == KERNEL_MIX_AIC_TASK_TYPE || taskType == KERNEL_MIX_AIV_TASK_TYPE) ? 0 : DEFAULT_CONTEXT_ID;
554 // 使用父节点的id作为connection_id，主要是为了将record_event的api与task_track关联起来598 // 使用父节点的id作为connection_id，主要是为了将record_event的api与task_track关联起来
555- auto task = GenHostTask(track, modelApi, nullptr, contextId,599+ std::shared_ptr<Operator> op;
556- track->data.runtimeTrack.taskType, node->parent->event->id);600+ std::shared_ptr<OpDesc> desc;
601+ MAKE_SHARED_RETURN_VALUE(desc, OpDesc, {});
602+ desc->runtimeTrackDesc = track;
603+ MAKE_SHARED_RETURN_VALUE(op, Operator, {}, desc, track->data.runtimeTrack.kernelName, OpType::OPTYPE_RESERVED);
604+ auto task = GenHostTask(track, modelApi, op, contextId, track->data.runtimeTrack.taskType, node->parent->event->id);
557 return task;605 return task;
558}606}
559 607 
560ComputeOpDescs TreeAnalyzer::GetComputeOpDescs(const std::shared_ptr<TreeNode> &nodeNode, bool useCtxId)608ComputeOpDescs TreeAnalyzer::GetComputeOpDescs(const std::shared_ptr<TreeNode> &nodeNode, bool useCtxId)
561{609{
562 ComputeOpDescs opDescs;610 ComputeOpDescs opDescs;
563- if (!nodeNode) {611+ if (!nodeNode)
612+ {
564 ERROR("nodeNode is nullptr, threadId = %", threadId_);613 ERROR("nodeNode is nullptr, threadId = %", threadId_);
565 return opDescs;614 return opDescs;
566 }615 }
567- for (const auto &record: nodeNode->records) {616+ for (const auto &record : nodeNode->records)
568- if (record->info.type == EventType::EVENT_TYPE_NODE_BASIC_INFO) {617+ {
618+ if (record->info.type == EventType::EVENT_TYPE_NODE_BASIC_INFO)
619+ {
569 std::shared_ptr<MsprofCompactInfo> trace;620 std::shared_ptr<MsprofCompactInfo> trace;
570 MAKE_SHARED_RETURN_VALUE(trace, MsprofCompactInfo, opDescs, *(record->compactPtr));621 MAKE_SHARED_RETURN_VALUE(trace, MsprofCompactInfo, opDescs, *(record->compactPtr));
571- UpdateComputeOpDescs<MsprofCompactInfo, &OpDesc::nodeDesc>(opDescs, trace,622+ UpdateComputeOpDescs<MsprofCompactInfo, &OpDesc::nodeDesc>(opDescs, trace, trace->data.nodeBasicInfo.opName,
572- trace->data.nodeBasicInfo.opName,
573 nodeNode->event->id);623 nodeNode->event->id);
574- } else if (record->info.type == EventType::EVENT_TYPE_NODE_ATTR_INFO) {624+ }
625+ else if (record->info.type == EventType::EVENT_TYPE_NODE_ATTR_INFO)
626+ {
575 std::shared_ptr<MsprofCompactInfo> trace;627 std::shared_ptr<MsprofCompactInfo> trace;
576 MAKE_SHARED_RETURN_VALUE(trace, MsprofCompactInfo, opDescs, *(record->compactPtr));628 MAKE_SHARED_RETURN_VALUE(trace, MsprofCompactInfo, opDescs, *(record->compactPtr));
577- UpdateComputeOpDescs<MsprofCompactInfo, &OpDesc::nodeAttr>(opDescs, trace,629+ UpdateComputeOpDescs<MsprofCompactInfo, &OpDesc::nodeAttr>(opDescs, trace, trace->data.nodeAttrInfo.opName,
578- trace->data.nodeAttrInfo.opName,
579 nodeNode->event->id);630 nodeNode->event->id);
580- } else if (record->info.type == EventType::EVENT_TYPE_TENSOR_INFO) {631+ }
632+ else if (record->info.type == EventType::EVENT_TYPE_TENSOR_INFO)
633+ {
581 std::shared_ptr<ConcatTensorInfo> trace;634 std::shared_ptr<ConcatTensorInfo> trace;
582 MAKE_SHARED_RETURN_VALUE(trace, ConcatTensorInfo, opDescs, *(record->tensorPtr));635 MAKE_SHARED_RETURN_VALUE(trace, ConcatTensorInfo, opDescs, *(record->tensorPtr));
583- UpdateComputeOpDescs<ConcatTensorInfo, &OpDesc::tensorDesc>(opDescs, trace,636+ UpdateComputeOpDescs<ConcatTensorInfo, &OpDesc::tensorDesc>(opDescs, trace, trace->opName,
584- trace->opName,
585 nodeNode->event->id);637 nodeNode->event->id);
586- } else if (record->info.type == EventType::EVENT_TYPE_CONTEXT_ID) {638+ }
587- if (!useCtxId) {639+ else if (record->info.type == EventType::EVENT_TYPE_CONTEXT_ID)
640+ {
641+ if (!useCtxId)
642+ {
588 continue;643 continue;
589 }644 }
590 std::shared_ptr<MsprofAdditionalInfo> trace;645 std::shared_ptr<MsprofAdditionalInfo> trace;
591 MAKE_SHARED_RETURN_VALUE(trace, MsprofAdditionalInfo, opDescs, *(record->additionPtr));646 MAKE_SHARED_RETURN_VALUE(trace, MsprofAdditionalInfo, opDescs, *(record->additionPtr));
592 auto ctxIdNode = ReinterpretConvert<MsprofContextIdInfo *>(trace->data);647 auto ctxIdNode = ReinterpretConvert<MsprofContextIdInfo *>(trace->data);
593- UpdateComputeOpDescs<MsprofAdditionalInfo, &OpDesc::ctxId>(opDescs, trace,648+ UpdateComputeOpDescs<MsprofAdditionalInfo, &OpDesc::ctxId>(opDescs, trace, ctxIdNode->opName,
594- ctxIdNode->opName,
595 nodeNode->event->id);649 nodeNode->event->id);
596- } else if (record->info.type == EventType::EVENT_TYPE_HCCL_OP_INFO) {650+ }
651+ else if (record->info.type == EventType::EVENT_TYPE_HCCL_OP_INFO)
652+ {
597 continue;653 continue;
598- } else {654+ }
599- ERROR("Unsupported additional type = %, threadId = %",655+ else
600- static_cast<uint16_t>(record->info.type), threadId_);656+ {
657+ ERROR("Unsupported additional type = %, threadId = %", static_cast<uint16_t>(record->info.type), threadId_);
601 }658 }
602 }659 }
603- if (!GetNodeRecordsByType(nodeNode, EventType::EVENT_TYPE_CONTEXT_ID).empty()) {660+ if (!GetNodeRecordsByType(nodeNode, EventType::EVENT_TYPE_CONTEXT_ID).empty())
661+ {
604 std::shared_ptr<OpDesc> desc;662 std::shared_ptr<OpDesc> desc;
605 MAKE_SHARED_RETURN_VALUE(desc, OpDesc, {});663 MAKE_SHARED_RETURN_VALUE(desc, OpDesc, {});
606 std::shared_ptr<Operator> op;664 std::shared_ptr<Operator> op;
@@ -620,19 +678,26 @@ HCCLSmallOpDescs TreeAnalyzer::GetHcclSmallOpDescs(const std::shared_ptr<TreeNod
620 auto hcclApi = hcclNode->event->apiPtr;678 auto hcclApi = hcclNode->event->apiPtr;
621 auto isMaster = TypeData::GetInstance().Get(hcclApi->level, hcclApi->type) == "master" ? 1 : 0;679 auto isMaster = TypeData::GetInstance().Get(hcclApi->level, hcclApi->type) == "master" ? 1 : 0;
622 680 
623- if (!ctxIdRecords.empty()) {681+ if (!ctxIdRecords.empty())
682+ {
624 // ctxId 存在，先根据其生成descs, 再补充hccl info683 // ctxId 存在，先根据其生成descs, 再补充hccl info
625 auto ret = UpdateHcclSmallOpDescs(opDescs, ctxIdRecords, hcclInfoRecords, isMaster);684 auto ret = UpdateHcclSmallOpDescs(opDescs, ctxIdRecords, hcclInfoRecords, isMaster);
626- if (!ret) {685+ if (!ret)
686+ {
627 ERROR("Update hccl small op descs failed by ctxId and hcclInfo, threadId = %", threadId_);687 ERROR("Update hccl small op descs failed by ctxId and hcclInfo, threadId = %", threadId_);
628 }688 }
629- } else if (!hcclInfoRecords.empty()) {689+ }
690+ else if (!hcclInfoRecords.empty())
691+ {
630 // 根据hccl info 生成692 // 根据hccl info 生成
631 auto ret = UpdateHcclSmallOpDescs(opDescs, hcclInfoRecords, isMaster);693 auto ret = UpdateHcclSmallOpDescs(opDescs, hcclInfoRecords, isMaster);
632- if (!ret) {694+ if (!ret)
695+ {
633 ERROR("Update hccl small op descs failed by hcclInfo, threadId = %", threadId_);696 ERROR("Update hccl small op descs failed by hcclInfo, threadId = %", threadId_);
634 }697 }
635- } else {698+ }
699+ else
700+ {
636 std::shared_ptr<HcclSmallOpDesc> desc;701 std::shared_ptr<HcclSmallOpDesc> desc;
637 MAKE_SHARED_RETURN_VALUE(desc, HcclSmallOpDesc, opDescs, DEFAULT_CONTEXT_ID, isMaster, nullptr);702 MAKE_SHARED_RETURN_VALUE(desc, HcclSmallOpDesc, opDescs, DEFAULT_CONTEXT_ID, isMaster, nullptr);
638 std::shared_ptr<Operator> op;703 std::shared_ptr<Operator> op;
@@ -645,18 +710,20 @@ HCCLSmallOpDescs TreeAnalyzer::GetHcclSmallOpDescs(const std::shared_ptr<TreeNod
645 710 
646bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,711bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,
647 const std::vector<std::shared_ptr<Event>> &ctxIdRecords,712 const std::vector<std::shared_ptr<Event>> &ctxIdRecords,
648- const std::vector<std::shared_ptr<Event>> &hcclInfoRecords,713+ const std::vector<std::shared_ptr<Event>> &hcclInfoRecords, uint8_t isMaster)
649- uint8_t isMaster)
650{714{
651 // 根据ctxId生成descs715 // 根据ctxId生成descs
652- for (const auto &record: ctxIdRecords) {716+ for (const auto &record : ctxIdRecords)
717+ {
653 auto trace = record->additionPtr;718 auto trace = record->additionPtr;
654 auto ctxIdTrace = ReinterpretConvert<MsprofContextIdInfo *>(trace->data);719 auto ctxIdTrace = ReinterpretConvert<MsprofContextIdInfo *>(trace->data);
655- if (ctxIdTrace->ctxIdNum != VALID_CTXID_NUM) {720+ if (ctxIdTrace->ctxIdNum != VALID_CTXID_NUM)
721+ {
656 ERROR("Expect ctxIdNum is 2, but get ctxIdNum is %, threadId = %", ctxIdTrace->ctxIdNum, threadId_);722 ERROR("Expect ctxIdNum is 2, but get ctxIdNum is %, threadId = %", ctxIdTrace->ctxIdNum, threadId_);
657 return false;723 return false;
658 }724 }
659- for (uint32_t id = ctxIdTrace->ctxIds[0]; id <= ctxIdTrace->ctxIds[1]; ++id) {725+ for (uint32_t id = ctxIdTrace->ctxIds[0]; id <= ctxIdTrace->ctxIds[1]; ++id)
726+ {
660 std::shared_ptr<HcclSmallOpDesc> desc;727 std::shared_ptr<HcclSmallOpDesc> desc;
661 MAKE_SHARED_RETURN_VALUE(desc, HcclSmallOpDesc, false, DEFAULT_CONTEXT_ID, isMaster, nullptr);728 MAKE_SHARED_RETURN_VALUE(desc, HcclSmallOpDesc, false, DEFAULT_CONTEXT_ID, isMaster, nullptr);
662 desc->ctxId = id;729 desc->ctxId = id;
@@ -665,7 +732,8 @@ bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,
665 descs.insert({id, op});732 descs.insert({id, op});
666 }733 }
667 }734 }
668- if (!ctxIdRecords.empty()) {735+ if (!ctxIdRecords.empty())
736+ {
669 std::shared_ptr<HcclSmallOpDesc> desc;737 std::shared_ptr<HcclSmallOpDesc> desc;
670 MAKE_SHARED_RETURN_VALUE(desc, HcclSmallOpDesc, false, DEFAULT_CONTEXT_ID, isMaster, nullptr);738 MAKE_SHARED_RETURN_VALUE(desc, HcclSmallOpDesc, false, DEFAULT_CONTEXT_ID, isMaster, nullptr);
671 std::shared_ptr<Operator> op;739 std::shared_ptr<Operator> op;
@@ -675,14 +743,18 @@ bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,
675 }743 }
676 744 
677 // HcclInfo更新745 // HcclInfo更新
678- for (const auto &record: hcclInfoRecords) {746+ for (const auto &record : hcclInfoRecords)
747+ {
679 auto trace = record->additionPtr;748 auto trace = record->additionPtr;
680 auto hcclTrace = ReinterpretConvert<MsprofHcclInfo *>(trace->data);749 auto hcclTrace = ReinterpretConvert<MsprofHcclInfo *>(trace->data);
681 auto key = hcclTrace->ctxID;750 auto key = hcclTrace->ctxID;
682- if (descs.find(key) != descs.end()) {751+ if (descs.find(key) != descs.end())
752+ {
683 auto hcclPtr = descs[key]->hcclSmallOpDesc;753 auto hcclPtr = descs[key]->hcclSmallOpDesc;
684 hcclPtr->hcclInfo = trace;754 hcclPtr->hcclInfo = trace;
685- } else {755+ }
756+ else
757+ {
686 ERROR("Can not find ctxId : % in descs, timestamp = %, threadId = %", key, trace->timeStamp, threadId_);758 ERROR("Can not find ctxId : % in descs, timestamp = %, threadId = %", key, trace->timeStamp, threadId_);
687 }759 }
688 }760 }
@@ -690,10 +762,10 @@ bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,
690}762}
691 763 
692bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,764bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,
693- const std::vector<std::shared_ptr<Event>> &hcclInfoRecords,765+ const std::vector<std::shared_ptr<Event>> &hcclInfoRecords, uint8_t isMaster)
694- uint8_t isMaster)
695{766{
696- for (const auto &record: hcclInfoRecords) {767+ for (const auto &record : hcclInfoRecords)
768+ {
697 auto trace = record->additionPtr;769 auto trace = record->additionPtr;
698 auto hcclTrace = ReinterpretConvert<MsprofHcclInfo *>(trace->data);770 auto hcclTrace = ReinterpretConvert<MsprofHcclInfo *>(trace->data);
699 auto key = hcclTrace->ctxID;771 auto key = hcclTrace->ctxID;
@@ -711,6 +783,6 @@ bool TreeAnalyzer::UpdateHcclSmallOpDescs(HCCLSmallOpDescs &descs,
711 return true;783 return true;
712}784}
713 785 
714-} // namespace Cann786+} // namespace Cann
715-} // namespace Association787+} // namespace Domain
716-} // namespace Analysis788+} // namespace Analysis
Manalysis/csrc/domain/services/persistence/host/cann_trace_db_dumper.cpp+86-26
@@ -56,9 +56,13 @@ const int32_t INVALID_VALUE = -1;
56const std::string NA = "N/A";56const std::string NA = "N/A";
57const uint32_t INPUT_FORMAT_INDEX = 0;57const uint32_t INPUT_FORMAT_INDEX = 0;
58const uint32_t OUTPUT_FORMAT_INDEX = 1;58const uint32_t OUTPUT_FORMAT_INDEX = 1;
59+const std::string KERNEL_MIX_AIC_TASK_TYPE = "KERNEL_MIX_AIC";
60+const std::string KERNEL_MIX_AIV_TASK_TYPE = "KERNEL_MIX_AIV";
61+const std::string KERNEL_SIMT_TASK_TYPE = "KERNEL_SIMT";
59const std::unordered_map<std::string, uint32_t> RtsTaskTypeMap = {62const std::unordered_map<std::string, uint32_t> RtsTaskTypeMap = {
60- {"KERNEL_AICORE", 0}, {"KERNEL_AICPU", 1}, {"KERNEL_AIVEC", 2}, {"KERNEL_MIX_AIC", 4}, {"KERNEL_MIX_AIV", 5},63+ {"KERNEL_AICORE", 0}, {"KERNEL_AICPU", 1}, {"KERNEL_AIVEC", 2},
61-};64+ {"KERNEL_MIX_AIC", 4}, {"KERNEL_MIX_AIV", 5}, {"KERNEL_SIMT", 2},
65+}; // KERNEL_SIMT类型算子是aiv类型的
62 66 
63std::string TransTaskTypeFromRtsToGe(uint64_t rtsTaskType)67std::string TransTaskTypeFromRtsToGe(uint64_t rtsTaskType)
64{68{
@@ -310,6 +314,10 @@ void CANNTraceDBDumper::AddTensorShapeInfo(const std::shared_ptr<ConcatTensorInf
310 auto mixBlockNum = blockNum * (nodeBasicInfo.blockNum >> 16);314 auto mixBlockNum = blockNum * (nodeBasicInfo.blockNum >> 16);
311 auto opFlag = nodeBasicInfo.opFlag ? "YES" : "NO";315 auto opFlag = nodeBasicInfo.opFlag ? "YES" : "NO";
312 auto opState = std::to_string(nodeBasicInfo.opState);316 auto opState = std::to_string(nodeBasicInfo.opState);
317+ auto runtimeTrackDesc = desc->runtimeTrackDesc;
318+ std::string gridDim = NA;
319+ std::string blockDim = NA;
320+ ProcessRuntimeTrackInfo(runtimeTrackDesc, blockNum, mixBlockNum, gridDim, blockDim);
313 auto inputFormatStr = inputFormat.empty() ? NA : Utils::Join(inputFormat, ";");321 auto inputFormatStr = inputFormat.empty() ? NA : Utils::Join(inputFormat, ";");
314 auto inputDataTypeStr = inputDataType.empty() ? NA : Utils::Join(inputDataType, ";");322 auto inputDataTypeStr = inputDataType.empty() ? NA : Utils::Join(inputDataType, ";");
315 auto inputShapeStr = inputShape.empty() ? NA : Utils::AddQuotation(Utils::Join(inputShape, ";"));323 auto inputShapeStr = inputShape.empty() ? NA : Utils::AddQuotation(Utils::Join(inputShape, ";"));
@@ -321,7 +329,7 @@ void CANNTraceDBDumper::AddTensorShapeInfo(const std::shared_ptr<ConcatTensorInf
321 mixBlockNum, opState, NumberMapping::Get(MappingType::GE_TASK_TYPE, nodeBasicInfo.taskType),329 mixBlockNum, opState, NumberMapping::Get(MappingType::GE_TASK_TYPE, nodeBasicInfo.taskType),
322 HashData::GetInstance().Get(nodeBasicInfo.opType), task->requestId, task->thread_id, task->timeStamp,330 HashData::GetInstance().Get(nodeBasicInfo.opType), task->requestId, task->thread_id, task->timeStamp,
323 task->batchId, tensorNum, inputFormatStr, inputDataTypeStr, inputShapeStr, outputFormatStr, outputDataTypeStr,331 task->batchId, tensorNum, inputFormatStr, inputDataTypeStr, inputShapeStr, outputFormatStr, outputDataTypeStr,
324- outputShapeStr, task->deviceId, task->contextId, opFlag, hashId);332+ outputShapeStr, task->deviceId, task->contextId, opFlag, hashId, gridDim, blockDim);
325}333}
326 334 
327std::string CANNTraceDBDumper::GetFormat(uint32_t oriFormat)335std::string CANNTraceDBDumper::GetFormat(uint32_t oriFormat)
@@ -369,36 +377,83 @@ void CANNTraceDBDumper::AddTaskInfoForOnlyTaskTrack(const std::shared_ptr<HostTa
369 opType = isLevel0 ? NA : kernelName;377 opType = isLevel0 ? NA : kernelName;
370 opName = kernelName;378 opName = kernelName;
371 }379 }
380+ std::string gridDim = NA;
381+ std::string blockDim = NA;
382+ auto runtimeTrackDesc = task->op->opDesc->runtimeTrackDesc;
372 if (isLevel0)383 if (isLevel0)
373 {384 {
374- data.emplace_back(info.modelId, opName, task->streamId, task->taskId, 0, 0, NA, taskType, opType,385+ uint32_t blockNum = 0;
375- task->requestId, task->thread_id, task->timeStamp, task->batchId, 0, NA, NA, NA, NA, NA, NA,386+ uint32_t mixBlockNum = 0;
376- task->deviceId, task->contextId, opFlag, info.hashId);387+ ProcessRuntimeTrackInfo(runtimeTrackDesc, blockNum, mixBlockNum, gridDim, blockDim);
388+ data.emplace_back(info.modelId, opName, task->streamId, task->taskId, blockNum, mixBlockNum, NA, taskType,
389+ opType, task->requestId, task->thread_id, task->timeStamp, task->batchId, 0, NA, NA, NA, NA,
390+ NA, NA, task->deviceId, task->contextId, opFlag, info.hashId, NA, NA);
377 }391 }
378 else392 else
379 {393 {
380- data.emplace_back(info.modelId, opName, task->streamId, task->taskId, info.blockNum, info.mixBlockNum,394+ uint32_t blockNum = info.blockNum;
381- info.isDynamic, taskType, opType, task->requestId, task->thread_id, task->timeStamp,395+ uint32_t mixBlockNum = info.mixBlockNum;
382- task->batchId, info.tensorNum, info.inputFormats, info.inputDataTypes, info.inputShapes,396+ ProcessRuntimeTrackInfo(runtimeTrackDesc, blockNum, mixBlockNum, gridDim, blockDim);
383- info.outputFormats, info.outputDataTypes, info.outputShapes, task->deviceId, task->contextId,397+ data.emplace_back(info.modelId, opName, task->streamId, task->taskId, blockNum, mixBlockNum, info.isDynamic,
384- opFlag, info.hashId);398+ taskType, opType, task->requestId, task->thread_id, task->timeStamp, task->batchId,
399+ info.tensorNum, info.inputFormats, info.inputDataTypes, info.inputShapes, info.outputFormats,
400+ info.outputDataTypes, info.outputShapes, task->deviceId, task->contextId, opFlag, info.hashId,
401+ gridDim, blockDim);
402+ }
403+}
404+ 
405+void CANNTraceDBDumper::ProcessRuntimeTrackInfo(const std::shared_ptr<MsprofCompactInfo> &runtimeTrack,
406+ uint32_t &blockNum, uint32_t &mixBlockNum, std::string &gridDim,
407+ std::string &blockDim)
408+{
409+ if (!runtimeTrack || runtimeTrack->dataLen != MSPROF_COMPACT_INFO_DATA_LENGTH)
410+ {
411+ return;
Wangang Yu
Wangang YuWangang Yu5月27日

[review] 问题： 当runtimeTrack为空或dataLen不匹配时，直接静默返回，没有任何日志记录 影响：数据解析失败时无法感知，问题被掩盖，无法快速定位是数据为空还是长度不匹配导致的问题 建议： 增加相关日志

likedislike
xfeng
xfeng
6月1日 评论：
412+ }
413+ auto taskType = TypeData::GetInstance().Get(MSPROF_REPORT_RUNTIME_LEVEL, runtimeTrack->data.runtimeTrack.taskType);
414+ if (taskType == KERNEL_SIMT_TASK_TYPE)
W
Wwangzixuan6月1日

[review] 这里的算子类型是rts提供的Kernel开头的算子类型风格。但是当前profiling数据提供通常不对外暴露rts风格的taskType。建议和其他类型一样，保持geTaskType风格。可参考onlyTaskTrack流程中的TransTaskTypeFromRtsToGe(task->taskType);方法

likedislike
xfeng
xfeng
6月1日 评论：
415+ {
416+ auto &simtInfo = runtimeTrack->data.runtimeTrack.extInfo.simtKernelInfo;
417+ gridDim =
418+ Utils::Join(std::vector<std::string>{std::to_string(simtInfo.gridDim.x), std::to_string(simtInfo.gridDim.y),
419+ std::to_string(simtInfo.gridDim.z)},
420+ ",");
421+ blockDim = Utils::Join(
422+ std::vector<std::string>{std::to_string(simtInfo.blockDim.x), std::to_string(simtInfo.blockDim.y),
423+ std::to_string(simtInfo.blockDim.z)},
424+ ",");
425+ }
426+ else
427+ {
428+ auto ratio = runtimeTrack->data.runtimeTrack.extInfo.kernelInfo.ratio;
429+ auto numBlocks = runtimeTrack->data.runtimeTrack.extInfo.kernelInfo.numBlocks;
430+ blockNum = numBlocks;
431+ mixBlockNum = numBlocks * ratio;
385 }432 }
386}433}
387 434 
388void CANNTraceDBDumper::AddTaskInfo(const std::shared_ptr<HostTask> &task, TaskInfoData &data, bool isLevel0)435void CANNTraceDBDumper::AddTaskInfo(const std::shared_ptr<HostTask> &task, TaskInfoData &data, bool isLevel0)
389{436{
390- if (!task->op)437+ if (!task->op || !task->op->opDesc)
438+ {
439+ return;
440+ }
441+ if (task->op->type == OpType::OPTYPE_RESERVED)
W
Wwangzixuan6月1日

[review] 这里逻辑被修改变更了。无nodeBasicInfo等情况下需要走onlyTaskTrack流程。这里的task->op->type == OpType::OPTYPE_RESERVED判定条件不能和原本流程对齐。建议明确相关流程，保持前后兼容统一。

likedislike
xfeng
xfeng
6月1日 评论：
391 {442 {
392 AddTaskInfoForOnlyTaskTrack(task, data, isLevel0);443 AddTaskInfoForOnlyTaskTrack(task, data, isLevel0);
393 return;444 return;
394 }445 }
395- 
396 if (isLevel0)446 if (isLevel0)
397 {447 {
448+ uint32_t blockNum = 0;
449+ uint32_t mixBlockNum = 0;
450+ std::string gridDim = NA;
451+ std::string blockDim = NA;
398 auto name = HashData::GetInstance().Get(task->op->name);452 auto name = HashData::GetInstance().Get(task->op->name);
399- data.emplace_back(task->modelId, name, task->streamId, task->taskId, 0, 0, NA, NA, NA, task->requestId,453+ ProcessRuntimeTrackInfo(task->op->opDesc->runtimeTrackDesc, blockNum, mixBlockNum, gridDim, blockDim);
400- task->thread_id, task->timeStamp, task->batchId, 0, NA, NA, NA, NA, NA, NA, task->deviceId,454+ data.emplace_back(task->modelId, name, task->streamId, task->taskId, blockNum, mixBlockNum, NA, NA, NA,
401- task->contextId, NA, NA);455+ task->requestId, task->thread_id, task->timeStamp, task->batchId, 0, NA, NA, NA, NA, NA, NA,
456+ task->deviceId, task->contextId, NA, NA, gridDim, blockDim);
402 return;457 return;
403 }458 }
404 459 
@@ -411,19 +466,24 @@ void CANNTraceDBDumper::AddTaskInfo(const std::shared_ptr<HostTask> &task, TaskI
411 auto node = desc->nodeDesc;466 auto node = desc->nodeDesc;
412 auto attr = desc->nodeAttr;467 auto attr = desc->nodeAttr;
413 auto nodeBasicInfo = node->data.nodeBasicInfo;468 auto nodeBasicInfo = node->data.nodeBasicInfo;
414- auto hashId = attr ? std::to_string(attr->data.nodeAttrInfo.hashId) : NA;
415- auto blockNum = nodeBasicInfo.blockNum & 0xffff;
416- auto mixBlockNum = blockNum * (nodeBasicInfo.blockNum >> 16);
417 auto tensorDesc = desc->tensorDesc;469 auto tensorDesc = desc->tensorDesc;
418- auto opFlag = nodeBasicInfo.opFlag ? "YES" : "NO";
419- auto opState = std::to_string(nodeBasicInfo.opState);
420 if (!tensorDesc)470 if (!tensorDesc)
421 {471 {
422- data.emplace_back(472+ auto hashId = attr ? std::to_string(attr->data.nodeAttrInfo.hashId) : NA;
423- task->modelId, HashData::GetInstance().Get(nodeBasicInfo.opName), task->streamId, task->taskId, blockNum,473+ auto blockNum = nodeBasicInfo.blockNum & 0xffff;
424- mixBlockNum, opState, NumberMapping::Get(MappingType::GE_TASK_TYPE, nodeBasicInfo.taskType),474+ auto mixBlockNum = blockNum * (nodeBasicInfo.blockNum >> 16);
425- HashData::GetInstance().Get(nodeBasicInfo.opType), task->requestId, task->thread_id, task->timeStamp,475+ auto opFlag = nodeBasicInfo.opFlag ? "YES" : "NO";
426- task->batchId, 0, NA, NA, NA, NA, NA, NA, task->deviceId, task->contextId, opFlag, hashId);476+ auto opState = std::to_string(nodeBasicInfo.opState);
477+ auto runtimeTrackDesc = desc->runtimeTrackDesc;
478+ std::string gridDim = NA;
479+ std::string blockDim = NA;
480+ ProcessRuntimeTrackInfo(runtimeTrackDesc, blockNum, mixBlockNum, gridDim, blockDim);
481+ data.emplace_back(task->modelId, HashData::GetInstance().Get(nodeBasicInfo.opName), task->streamId,
482+ task->taskId, blockNum, mixBlockNum, opState,
483+ NumberMapping::Get(MappingType::GE_TASK_TYPE, nodeBasicInfo.taskType),
484+ HashData::GetInstance().Get(nodeBasicInfo.opType), task->requestId, task->thread_id,
485+ task->timeStamp, task->batchId, 0, NA, NA, NA, NA, NA, NA, task->deviceId, task->contextId,
486+ opFlag, hashId, gridDim, blockDim);
427 return;487 return;
428 }488 }
429 AddTensorShapeInfo(tensorDesc, nodeBasicInfo, data, task);489 AddTensorShapeInfo(tensorDesc, nodeBasicInfo, data, task);
Manalysis/csrc/domain/services/persistence/host/cann_trace_db_dumper.h+28-22
@@ -17,39 +17,44 @@
17#ifndef ANALYSIS_VIEWER_DATABASE_DRAFTS_CANN_DB_DUMPER_H17#ifndef ANALYSIS_VIEWER_DATABASE_DRAFTS_CANN_DB_DUMPER_H
18#define ANALYSIS_VIEWER_DATABASE_DRAFTS_CANN_DB_DUMPER_H18#define ANALYSIS_VIEWER_DATABASE_DRAFTS_CANN_DB_DUMPER_H
19 19 
20+#include <mutex>
21+#include <thread>
20#include <utility>22#include <utility>
21#include <vector>23#include <vector>
22-#include <thread>
23-#include <mutex>
24 24 
25-#include "analysis/csrc/domain/services/association/cann/include/tree_analyzer.h"
26#include "analysis/csrc/domain/entities/tree/include/tree.h"25#include "analysis/csrc/domain/entities/tree/include/tree.h"
27-#include "analysis/csrc/infrastructure/utils/utils.h"26+#include "analysis/csrc/domain/services/association/cann/include/tree_analyzer.h"
28#include "analysis/csrc/infrastructure/db/include/database.h"27#include "analysis/csrc/infrastructure/db/include/database.h"
29#include "analysis/csrc/infrastructure/db/include/db_runner.h"28#include "analysis/csrc/infrastructure/db/include/db_runner.h"
29+#include "analysis/csrc/infrastructure/utils/utils.h"
30 30 
31-namespace Analysis {31+namespace Analysis
32-namespace Domain {32+{
33+namespace Domain
34+{
33 35 
34// 供HostTraceWorker调用， 传入Analyzer对象，将所有数据落盘36// 供HostTraceWorker调用， 传入Analyzer对象，将所有数据落盘
35-class CANNTraceDBDumper {37+class CANNTraceDBDumper
36-using TreeAnalyzer = Analysis::Domain::Cann::TreeAnalyzer;38+{
37-using HostTask = Analysis::Domain::HostTask;39+ using TreeAnalyzer = Analysis::Domain::Cann::TreeAnalyzer;
38-using OpDesc = Analysis::Domain::OpDesc;40+ using HostTask = Analysis::Domain::HostTask;
39-using HostTasks = std::vector<std::shared_ptr<HostTask>>;41+ using OpDesc = Analysis::Domain::OpDesc;
40-using TaskInfoData = std::vector<std::tuple<uint32_t, std::string, uint32_t, uint32_t, uint32_t, uint32_t, std::string,42+ using HostTasks = std::vector<std::shared_ptr<HostTask>>;
41- std::string, std::string, int32_t, uint32_t, uint64_t, uint32_t, uint32_t, std::string, std::string,43+ using TaskInfoData = std::vector<
42- std::string, std::string, std::string, std::string, uint32_t, uint32_t, std::string, std::string>>;44+ std::tuple<uint32_t, std::string, uint32_t, uint32_t, uint32_t, uint32_t, std::string, std::string, std::string,
43-using HCCLBigOpDescs = Analysis::Domain::Cann::HCCLBigOpDescs;45+ int32_t, uint32_t, uint64_t, uint32_t, uint32_t, std::string, std::string, std::string, std::string,
44-using GeFusionOpInfos = Analysis::Domain::Cann::GeFusionOpInfos;46+ std::string, std::string, uint32_t, uint32_t, std::string, std::string, std::string, std::string>>;
45-public:47+ using HCCLBigOpDescs = Analysis::Domain::Cann::HCCLBigOpDescs;
48+ using GeFusionOpInfos = Analysis::Domain::Cann::GeFusionOpInfos;
49+ 
50+ public:
46 // 创建时传入host路径51 // 创建时传入host路径
47 explicit CANNTraceDBDumper(std::string hostFilePath);52 explicit CANNTraceDBDumper(std::string hostFilePath);
48 53 
49 // 提供DumpData方法，调用后执行落盘操作，成功返回True，失败返回False。54 // 提供DumpData方法，调用后执行落盘操作，成功返回True，失败返回False。
50 bool DumpData(TreeAnalyzer analyzer);55 bool DumpData(TreeAnalyzer analyzer);
51 56 
52-private:57+ private:
53 // 落盘HCCLOP58 // 落盘HCCLOP
54 void DumpHcclOps(const HCCLBigOpDescs &hcclOps);59 void DumpHcclOps(const HCCLBigOpDescs &hcclOps);
55 60 
@@ -70,13 +75,14 @@ private:
70 75 
71 void AddTensorShapeInfo(const std::shared_ptr<ConcatTensorInfo> &tensorDesc, MsprofNodeBasicInfo nodeBasicInfo,76 void AddTensorShapeInfo(const std::shared_ptr<ConcatTensorInfo> &tensorDesc, MsprofNodeBasicInfo nodeBasicInfo,
72 TaskInfoData &data, const std::shared_ptr<HostTask> &task);77 TaskInfoData &data, const std::shared_ptr<HostTask> &task);
78+ static void ProcessRuntimeTrackInfo(const std::shared_ptr<MsprofCompactInfo> &runtimeTrack, uint32_t &blockNum,
79+ uint32_t &mixBlockNum, std::string &gridDimStr, std::string &blockDimStr);
73 static std::string GetFormat(uint32_t oriFormat);80 static std::string GetFormat(uint32_t oriFormat);
74 const uint32_t poolSize_ = 4;81 const uint32_t poolSize_ = 4;
75 std::string hostFilePath_;82 std::string hostFilePath_;
76 std::atomic<bool> result_;83 std::atomic<bool> result_;
77};84};
78-} // Domain85+} // namespace Domain
79-} // Analysis86+} // namespace Analysis
80 87 
81- 88+#endif // ANALYSIS_VIEWER_DATABASE_DRAFTS_CANN_DB_DUMPER_H
82-#endif // ANALYSIS_VIEWER_DATABASE_DRAFTS_CANN_DB_DUMPER_H
Manalysis/csrc/infrastructure/db/database.cpp+2-1
@@ -43,7 +43,8 @@ const TableColumns TaskInfo = {
43 {"batch_id", SQL_INTEGER_TYPE}, {"tensor_num", SQL_INTEGER_TYPE}, {"input_formats", SQL_TEXT_TYPE},43 {"batch_id", SQL_INTEGER_TYPE}, {"tensor_num", SQL_INTEGER_TYPE}, {"input_formats", SQL_TEXT_TYPE},
44 {"input_data_types", SQL_TEXT_TYPE}, {"input_shapes", SQL_TEXT_TYPE}, {"output_formats", SQL_TEXT_TYPE},44 {"input_data_types", SQL_TEXT_TYPE}, {"input_shapes", SQL_TEXT_TYPE}, {"output_formats", SQL_TEXT_TYPE},
45 {"output_data_types", SQL_TEXT_TYPE}, {"output_shapes", SQL_TEXT_TYPE}, {"device_id", SQL_INTEGER_TYPE},45 {"output_data_types", SQL_TEXT_TYPE}, {"output_shapes", SQL_TEXT_TYPE}, {"device_id", SQL_INTEGER_TYPE},
46- {"context_id", SQL_INTEGER_TYPE}, {"op_flag", SQL_TEXT_TYPE}, {"hashid", SQL_TEXT_TYPE}};46+ {"context_id", SQL_INTEGER_TYPE}, {"op_flag", SQL_TEXT_TYPE}, {"hashid", SQL_TEXT_TYPE},
47+ {"grid_dim", SQL_TEXT_TYPE}, {"block_dim", SQL_TEXT_TYPE}};
47 48 
48const TableColumns StepInfo = {{"model_id", SQL_INTEGER_TYPE},49const TableColumns StepInfo = {{"model_id", SQL_INTEGER_TYPE},
49 {"thread_id", SQL_INTEGER_TYPE},50 {"thread_id", SQL_INTEGER_TYPE},
Manalysis/csrc/infrastructure/utils/prof_common.h+486-414
@@ -17,139 +17,154 @@
17#define MSPROFILER_PROF_COMMON_H_17#define MSPROFILER_PROF_COMMON_H_
18 18 
19#include <stdint.h>19#include <stdint.h>
20+ 
20#include <vector>21#include <vector>
21 22 
22#ifdef __cplusplus23#ifdef __cplusplus
23-extern "C" {24+extern "C"
24-#endif // __cplusplus25+{
26+#endif // __cplusplus
25 27 
26-#define MSPROF_DATA_HEAD_MAGIC_NUM 0x5a5a28+#define MSPROF_DATA_HEAD_MAGIC_NUM 0x5a5a
27#define MSPROF_EVENT_FLAG 0xFFFFFFFFFFFFFFFFULL29#define MSPROF_EVENT_FLAG 0xFFFFFFFFFFFFFFFFULL
28 30 
29-enum MsprofDataTag {31+ enum MsprofDataTag
30- MSPROF_ACL_DATA_TAG = 0, // acl data tag, range: 0~1932+ {
31- MSPROF_GE_DATA_TAG_MODEL_LOAD = 20, // ge data tag, range: 20~3933+ MSPROF_ACL_DATA_TAG = 0, // acl data tag, range: 0~19
32- MSPROF_GE_DATA_TAG_FUSION = 21,34+ MSPROF_GE_DATA_TAG_MODEL_LOAD = 20, // ge data tag, range: 20~39
33- MSPROF_GE_DATA_TAG_INFER = 22,35+ MSPROF_GE_DATA_TAG_FUSION = 21,
34- MSPROF_GE_DATA_TAG_TASK = 23,36+ MSPROF_GE_DATA_TAG_INFER = 22,
35- MSPROF_GE_DATA_TAG_TENSOR = 24,37+ MSPROF_GE_DATA_TAG_TASK = 23,
36- MSPROF_GE_DATA_TAG_STEP = 25,38+ MSPROF_GE_DATA_TAG_TENSOR = 24,
37- MSPROF_GE_DATA_TAG_ID_MAP = 26,39+ MSPROF_GE_DATA_TAG_STEP = 25,
38- MSPROF_GE_DATA_TAG_HOST_SCH = 27,40+ MSPROF_GE_DATA_TAG_ID_MAP = 26,
39- MSPROF_RUNTIME_DATA_TAG_API = 40, // runtime data tag, range: 40~5941+ MSPROF_GE_DATA_TAG_HOST_SCH = 27,
40- MSPROF_RUNTIME_DATA_TAG_TRACK = 41,42+ MSPROF_RUNTIME_DATA_TAG_API = 40, // runtime data tag, range: 40~59
41- MSPROF_AICPU_DATA_TAG = 60, // aicpu data tag, range: 60~7943+ MSPROF_RUNTIME_DATA_TAG_TRACK = 41,
42- MSPROF_AICPU_MODEL_TAG = 61,44+ MSPROF_AICPU_DATA_TAG = 60, // aicpu data tag, range: 60~79
43- MSPROF_HCCL_DATA_TAG = 80, // hccl data tag, range: 80~9945+ MSPROF_AICPU_MODEL_TAG = 61,
44- MSPROF_DP_DATA_TAG = 100, // dp data tag, range: 100~11946+ MSPROF_HCCL_DATA_TAG = 80, // hccl data tag, range: 80~99
45- MSPROF_MSPROFTX_DATA_TAG = 120, // msproftx data tag, range: 120~13947+ MSPROF_DP_DATA_TAG = 100, // dp data tag, range: 100~119
46- MSPROF_DATA_TAG_MAX = 65536, // data tag value type is uint16_t48+ MSPROF_MSPROFTX_DATA_TAG = 120, // msproftx data tag, range: 120~139
47-};49+ MSPROF_DATA_TAG_MAX = 65536, // data tag value type is uint16_t
50+ };
48 51 
49#define PATH_LEN_MAX 102352#define PATH_LEN_MAX 1023
50#define PARAM_LEN_MAX 409553#define PARAM_LEN_MAX 4095
51-struct MsprofCommandHandleParams {54+ struct MsprofCommandHandleParams
52- uint32_t pathLen;55+ {
53- uint32_t storageLimit; // MB56+ uint32_t pathLen;
54- uint32_t profDataLen;57+ uint32_t storageLimit; // MB
55- char path[PATH_LEN_MAX + 1];58+ uint32_t profDataLen;
56- char profData[PARAM_LEN_MAX + 1];59+ char path[PATH_LEN_MAX + 1];
57-};60+ char profData[PARAM_LEN_MAX + 1];
61+ };
58 62 
59/**63/**
60 * @brief profiling command info64 * @brief profiling command info
61 */65 */
62#define MSPROF_MAX_DEV_NUM 6466#define MSPROF_MAX_DEV_NUM 64
63-struct MsprofCommandHandle {67+ struct MsprofCommandHandle
64- uint64_t profSwitch;68+ {
65- uint64_t profSwitchHi;69+ uint64_t profSwitch;
66- uint32_t devNums;70+ uint64_t profSwitchHi;
67- uint32_t devIdList[MSPROF_MAX_DEV_NUM];71+ uint32_t devNums;
68- uint32_t modelId;72+ uint32_t devIdList[MSPROF_MAX_DEV_NUM];
69- uint32_t type;73+ uint32_t modelId;
70- uint32_t cacheFlag;74+ uint32_t type;
71- struct MsprofCommandHandleParams params;75+ uint32_t cacheFlag;
72-};76+ struct MsprofCommandHandleParams params;
77+ };
73 78 
74#define MSPROF_GE_TENSOR_DATA_SHAPE_LEN 879#define MSPROF_GE_TENSOR_DATA_SHAPE_LEN 8
75#define MSPROF_GE_TENSOR_DATA_NUM 580#define MSPROF_GE_TENSOR_DATA_NUM 5
76#define MSPROF_GE_FUSION_OP_NUM 881#define MSPROF_GE_FUSION_OP_NUM 8
77#define MSPROF_CTX_ID_MAX_NUM 5582#define MSPROF_CTX_ID_MAX_NUM 55
78#pragma pack(1)83#pragma pack(1)
79-struct MsprofNodeBasicInfo {84+ struct MsprofNodeBasicInfo
80- uint64_t opName;85+ {
81- uint32_t taskType;86+ uint64_t opName;
82- uint64_t opType;87+ uint32_t taskType;
83- uint32_t blockNum;88+ uint64_t opType;
84- uint32_t opFlag;89+ uint32_t blockNum;
85- uint8_t opState;90+ uint32_t opFlag;
86-};91+ uint8_t opState;
92+ };
87 93 
88-enum AttrType {94+ enum AttrType
89- OP_ATTR = 0,95+ {
90-};96+ OP_ATTR = 0,
97+ };
91 98 
92-struct MsprofAttrInfo {99+ struct MsprofAttrInfo
93- uint64_t opName;100+ {
94- uint32_t attrType;101+ uint64_t opName;
95- uint64_t hashId;102+ uint32_t attrType;
96-};103+ uint64_t hashId;
104+ };
97 105 
98-struct MsrofTensorData {106+ struct MsrofTensorData
99- uint32_t tensorType;107+ {
100- uint32_t format;108+ uint32_t tensorType;
101- uint32_t dataType;109+ uint32_t format;
102- uint32_t shape[MSPROF_GE_TENSOR_DATA_SHAPE_LEN];110+ uint32_t dataType;
103-};111+ uint32_t shape[MSPROF_GE_TENSOR_DATA_SHAPE_LEN];
112+ };
104 113 
105-struct MsprofTensorInfo {114+ struct MsprofTensorInfo
106- uint64_t opName;115+ {
107- uint32_t tensorNum;116+ uint64_t opName;
108- MsrofTensorData tensorData[MSPROF_GE_TENSOR_DATA_NUM];117+ uint32_t tensorNum;
109-};118+ MsrofTensorData tensorData[MSPROF_GE_TENSOR_DATA_NUM];
119+ };
110 120 
111-struct ProfFusionOpInfo {121+ struct ProfFusionOpInfo
112- uint64_t opName;122+ {
113- uint32_t fusionOpNum;123+ uint64_t opName;
114- uint64_t inputMemsize;124+ uint32_t fusionOpNum;
115- uint64_t outputMemsize;125+ uint64_t inputMemsize;
116- uint64_t weightMemSize;126+ uint64_t outputMemsize;
117- uint64_t workspaceMemSize;127+ uint64_t weightMemSize;
118- uint64_t totalMemSize;128+ uint64_t workspaceMemSize;
119- uint64_t fusionOpId[MSPROF_GE_FUSION_OP_NUM];129+ uint64_t totalMemSize;
120-};130+ uint64_t fusionOpId[MSPROF_GE_FUSION_OP_NUM];
131+ };
121 132 
122-struct MsprofContextIdInfo {133+ struct MsprofContextIdInfo
123- uint64_t opName;134+ {
124- uint32_t ctxIdNum;135+ uint64_t opName;
125- uint32_t ctxIds[MSPROF_CTX_ID_MAX_NUM];136+ uint32_t ctxIdNum;
126-};137+ uint32_t ctxIds[MSPROF_CTX_ID_MAX_NUM];
138+ };
127 139 
128-struct MsprofGraphIdInfo {140+ struct MsprofGraphIdInfo
129- uint64_t modelName;141+ {
130- uint32_t graphId;142+ uint64_t modelName;
131- uint32_t modelId;143+ uint32_t graphId;
132-};144+ uint32_t modelId;
145+ };
133 146 
134-struct MsprofMemoryInfo {147+ struct MsprofMemoryInfo
135- uint64_t addr;148+ {
136- int64_t size;149+ uint64_t addr;
137- uint64_t nodeId;150+ int64_t size;
138- uint64_t totalAllocateMemory;151+ uint64_t nodeId;
139- uint64_t totalReserveMemory;152+ uint64_t totalAllocateMemory;
140- uint32_t deviceId;153+ uint64_t totalReserveMemory;
141- uint32_t deviceType;154+ uint32_t deviceId;
142-};155+ uint32_t deviceType;
156+ };
143 157 
144-struct MsprofStaticOpMem {158+ struct MsprofStaticOpMem
145- int64_t size; // op memory size159+ {
146- uint64_t opName; // op name hash id160+ int64_t size; // op memory size
147- uint64_t lifeStart; // serial number of op memory used161+ uint64_t opName; // op name hash id
148- uint64_t lifeEnd; // serial number of op memory used162+ uint64_t lifeStart; // serial number of op memory used
149- uint64_t totalAllocateMemory; // static graph total allocate memory163+ uint64_t lifeEnd; // serial number of op memory used
150- uint64_t dynOpName; // 0: invalid， other： dynamic op name of root164+ uint64_t totalAllocateMemory; // static graph total allocate memory
151- uint32_t graphId; // multiple model165+ uint64_t dynOpName; // 0: invalid， other： dynamic op name of root
152-};166+ uint32_t graphId; // multiple model
167+ };
153 168 
154/**169/**
155 * @name MsprofStampInfo170 * @name MsprofStampInfo
@@ -157,40 +172,44 @@ struct MsprofStaticOpMem {
157 */172 */
158#define UIF_VALUE_LEN 2173#define UIF_VALUE_LEN 2
159#define MAX_MESSAGE_LEN 156174#define MAX_MESSAGE_LEN 156
160-struct MsprofStampInfo {175+ struct MsprofStampInfo
161- uint16_t magicNumber;176+ {
162- uint16_t dataTag;177+ uint16_t magicNumber;
163- uint32_t processId;178+ uint16_t dataTag;
164- uint32_t threadId;179+ uint32_t processId;
165- uint32_t category; // marker category180+ uint32_t threadId;
166- uint32_t eventType;181+ uint32_t category; // marker category
167- int32_t payloadType;182+ uint32_t eventType;
168- union PayloadValue {183+ int32_t payloadType;
169- uint64_t ullValue;184+ union PayloadValue
170- int64_t llValue;185+ {
171- double dValue;186+ uint64_t ullValue;
172- uint32_t uiValue[UIF_VALUE_LEN];187+ int64_t llValue;
173- int32_t iValue[UIF_VALUE_LEN];188+ double dValue;
174- float fValue[UIF_VALUE_LEN];189+ uint32_t uiValue[UIF_VALUE_LEN];
175- } payload; // payload info for marker190+ int32_t iValue[UIF_VALUE_LEN];
176- uint64_t startTime;191+ float fValue[UIF_VALUE_LEN];
177- uint64_t endTime;192+ } payload; // payload info for marker
178- uint64_t markId;193+ uint64_t startTime;
179- uint64_t domain;194+ uint64_t endTime;
180- int32_t messageType;195+ uint64_t markId;
181- char message[MAX_MESSAGE_LEN];196+ uint64_t domain;
182-};197+ int32_t messageType;
198+ char message[MAX_MESSAGE_LEN];
199+ };
183 200 
184-#define MSPROF_TX_VALUE_MAX_LEN 224 // 224 + 8 = 232: additional data len201+#define MSPROF_TX_VALUE_MAX_LEN 224 // 224 + 8 = 232: additional data len
185-struct MsprofTxInfo {202+ struct MsprofTxInfo
186- uint16_t infoType; // 0: Mark; 1: MarkEx203+ {
187- uint16_t res0;204+ uint16_t infoType; // 0: Mark; 1: MarkEx
188- uint32_t res1;205+ uint16_t res0;
189- union {206+ uint32_t res1;
190- struct MsprofStampInfo stampInfo;207+ union
191- uint8_t data[MSPROF_TX_VALUE_MAX_LEN];208+ {
192- } value;209+ struct MsprofStampInfo stampInfo;
193-};210+ uint8_t data[MSPROF_TX_VALUE_MAX_LEN];
211+ } value;
212+ };
194 213 
195#pragma pack()214#pragma pack()
196 215 
@@ -198,306 +217,359 @@ struct MsprofTxInfo {
198 * @brief struct of data reported by HCCL217 * @brief struct of data reported by HCCL
199 */218 */
200#pragma pack(4)219#pragma pack(4)
201-struct MsprofHcclInfo {220+ struct MsprofHcclInfo
202- uint64_t itemId;221+ {
203- uint64_t cclTag;222+ uint64_t itemId;
204- uint64_t groupName;223+ uint64_t cclTag;
205- uint32_t localRank;224+ uint64_t groupName;
206- uint32_t remoteRank;225+ uint32_t localRank;
207- uint32_t rankSize;226+ uint32_t remoteRank;
208- uint32_t workFlowMode;227+ uint32_t rankSize;
209- uint32_t planeID;228+ uint32_t workFlowMode;
210- uint32_t ctxID;229+ uint32_t planeID;
211- uint64_t notifyID;230+ uint32_t ctxID;
212- uint32_t stage;231+ uint64_t notifyID;
213- uint32_t role; // role {0: dst, 1:src}232+ uint32_t stage;
214- double durationEstimated;233+ uint32_t role; // role {0: dst, 1:src}
215- uint64_t srcAddr;234+ double durationEstimated;
216- uint64_t dstAddr;235+ uint64_t srcAddr;
217- uint64_t dataSize; // bytes236+ uint64_t dstAddr;
218- uint32_t opType; // {0: sum, 1: mul, 2: max, 3: min}237+ uint64_t dataSize; // bytes
219- uint32_t dataType; // data type {0: INT8, 1: INT16, 2: INT32, 3: FP16, 4:FP32, 5:INT64, 6:UINT64}238+ uint32_t opType; // {0: sum, 1: mul, 2: max, 3: min}
220- uint32_t linkType; // link type {0: 'OnChip', 1: 'HCCS', 2: 'PCIe', 3: 'RoCE', 4: 'SIO'}239+ uint32_t dataType; // data type {0: INT8, 1: INT16, 2: INT32, 3: FP16, 4:FP32, 5:INT64, 6:UINT64}
221- uint32_t transportType; // transport type {0: SDMA, 1: RDMA, 2:LOCAL}240+ uint32_t linkType; // link type {0: 'OnChip', 1: 'HCCS', 2: 'PCIe', 3: 'RoCE', 4: 'SIO'}
222- uint32_t rdmaType; // RDMA type {0: RDMASendNotify, 1:RDMASendPayload}241+ uint32_t transportType; // transport type {0: SDMA, 1: RDMA, 2:LOCAL}
223- uint32_t reserve2;242+ uint32_t rdmaType; // RDMA type {0: RDMASendNotify, 1:RDMASendPayload}
224-};243+ uint32_t reserve2;
244+ };
225 245 
226-const uint16_t MSPROF_MULTI_THREAD_MAX_NUM = 25;246+ const uint16_t MSPROF_MULTI_THREAD_MAX_NUM = 25;
227-struct MsprofMultiThread {247+ struct MsprofMultiThread
228- uint32_t threadNum;248+ {
229- uint32_t threadId[MSPROF_MULTI_THREAD_MAX_NUM];249+ uint32_t threadNum;
230-};250+ uint32_t threadId[MSPROF_MULTI_THREAD_MAX_NUM];
251+ };
231#pragma pack()252#pragma pack()
232 253 
233-/* Msprof report level */254+ /* Msprof report level */
234-const uint16_t MSPROF_REPORT_PYTORCH_LEVEL = 30000;255+ const uint16_t MSPROF_REPORT_PYTORCH_LEVEL = 30000;
235-const uint16_t MSPROF_REPORT_PTA_LEVEL = 25000;256+ const uint16_t MSPROF_REPORT_PTA_LEVEL = 25000;
236-const uint16_t MSPROF_REPORT_TX_LEVEL = 20500;257+ const uint16_t MSPROF_REPORT_TX_LEVEL = 20500;
237-const uint16_t MSPROF_REPORT_ACL_LEVEL = 20000;258+ const uint16_t MSPROF_REPORT_ACL_LEVEL = 20000;
238-const uint16_t MSPROF_REPORT_MODEL_LEVEL = 15000;259+ const uint16_t MSPROF_REPORT_MODEL_LEVEL = 15000;
239-const uint16_t MSPROF_REPORT_NODE_LEVEL = 10000;260+ const uint16_t MSPROF_REPORT_NODE_LEVEL = 10000;
240-const uint16_t MSPROF_REPORT_HCCL_NODE_LEVEL = 5500;261+ const uint16_t MSPROF_REPORT_HCCL_NODE_LEVEL = 5500;
241-const uint16_t MSPROF_REPORT_RUNTIME_LEVEL = 5000;262+ const uint16_t MSPROF_REPORT_RUNTIME_LEVEL = 5000;
242 263 
243-/* Msprof report type of pytorch(30000) level(proftx), offset: 0 */264+ /* Msprof report type of pytorch(30000) level(proftx), offset: 0 */
244-const uint32_t MSPROF_REPORT_PYTORCH_PROFTX_TYPE = 0;265+ const uint32_t MSPROF_REPORT_PYTORCH_PROFTX_TYPE = 0;
245-const uint32_t MSPROF_REPORT_PYTORCH_CATEGORY_DIC_TYPE = 1;266+ const uint32_t MSPROF_REPORT_PYTORCH_CATEGORY_DIC_TYPE = 1;
246-const uint32_t MSPROF_REPORT_PYTORCH_CALLSTACK_TYPE = 2;267+ const uint32_t MSPROF_REPORT_PYTORCH_CALLSTACK_TYPE = 2;
247-const uint32_t MSPROF_REPORT_PYTORCH_CANN_OP_TYPE = 3;268+ const uint32_t MSPROF_REPORT_PYTORCH_CANN_OP_TYPE = 3;
248-const uint32_t MSPROF_REPORT_PYTORCH_TORCH_OP_TYPE = 4;269+ const uint32_t MSPROF_REPORT_PYTORCH_TORCH_OP_TYPE = 4;
249-const uint32_t MSPROF_REPORT_PYTORCH_PIPELINE_TYPE = 5;270+ const uint32_t MSPROF_REPORT_PYTORCH_PIPELINE_TYPE = 5;
250 271 
251-/* Msprof report type of tx(20500) level, offset: 0x000000 */272+ /* Msprof report type of tx(20500) level, offset: 0x000000 */
252-const uint32_t MSPROF_REPORT_TX_BASE_TYPE = 0x000000U;273+ const uint32_t MSPROF_REPORT_TX_BASE_TYPE = 0x000000U;
253 274 
254-/* Msprof report type of acl(20000) level(acl), offset: 0x020000 */275+ /* Msprof report type of acl(20000) level(acl), offset: 0x020000 */
255-const uint32_t MSPROF_REPORT_ACL_OP_BASE_TYPE = 0x010000U;276+ const uint32_t MSPROF_REPORT_ACL_OP_BASE_TYPE = 0x010000U;
256-const uint32_t MSPROF_REPORT_ACL_MODEL_BASE_TYPE = 0x020000U;277+ const uint32_t MSPROF_REPORT_ACL_MODEL_BASE_TYPE = 0x020000U;
257-const uint32_t MSPROF_REPORT_ACL_RUNTIME_BASE_TYPE = 0x030000U;278+ const uint32_t MSPROF_REPORT_ACL_RUNTIME_BASE_TYPE = 0x030000U;
258-const uint32_t MSPROF_REPORT_ACL_OTHERS_BASE_TYPE = 0x040000U;279+ const uint32_t MSPROF_REPORT_ACL_OTHERS_BASE_TYPE = 0x040000U;
259 280 
260-/* Msprof report type of acl(20000) level(host api hccl), offset: 0x070000 */281+ /* Msprof report type of acl(20000) level(host api hccl), offset: 0x070000 */
261-const uint32_t MSPROF_REPORT_ACL_NN_BASE_TYPE = 0x050000U;282+ const uint32_t MSPROF_REPORT_ACL_NN_BASE_TYPE = 0x050000U;
262-const uint32_t MSPROF_REPORT_ACL_ASCENDC_TYPE = 0x060000U;283+ const uint32_t MSPROF_REPORT_ACL_ASCENDC_TYPE = 0x060000U;
263-const uint32_t MSPROF_REPORT_ACL_HOST_HCCL_BASE_TYPE = 0x070000U;284+ const uint32_t MSPROF_REPORT_ACL_HOST_HCCL_BASE_TYPE = 0x070000U;
264-const uint32_t MSPROF_REPORT_ACL_DVPP_BASE_TYPE = 0x090000U;285+ const uint32_t MSPROF_REPORT_ACL_DVPP_BASE_TYPE = 0x090000U;
265-const uint32_t MSPROF_REPORT_ACL_GRAPH_BASE_TYPE = 0x0A0000U;286+ const uint32_t MSPROF_REPORT_ACL_GRAPH_BASE_TYPE = 0x0A0000U;
266 287 
267-/* Msprof report type of model(15000) level, offset: 0x000000 */288+ /* Msprof report type of model(15000) level, offset: 0x000000 */
268-const uint32_t MSPROF_REPORT_MODEL_GRAPH_ID_MAP_TYPE = 0; /* type info: graph_id_map */289+ const uint32_t MSPROF_REPORT_MODEL_GRAPH_ID_MAP_TYPE = 0; /* type info: graph_id_map */
269-const uint32_t MSPROF_REPORT_MODEL_EXECUTE_TYPE = 1; /* type info: execute */290+ const uint32_t MSPROF_REPORT_MODEL_EXECUTE_TYPE = 1; /* type info: execute */
270-const uint32_t MSPROF_REPORT_MODEL_LOAD_TYPE = 2; /* type info: load */291+ const uint32_t MSPROF_REPORT_MODEL_LOAD_TYPE = 2; /* type info: load */
271-const uint32_t MSPROF_REPORT_MODEL_INPUT_COPY_TYPE = 3; /* type info: IntputCopy */292+ const uint32_t MSPROF_REPORT_MODEL_INPUT_COPY_TYPE = 3; /* type info: IntputCopy */
272-const uint32_t MSPROF_REPORT_MODEL_OUTPUT_COPY_TYPE = 4; /* type info: OutputCopy */293+ const uint32_t MSPROF_REPORT_MODEL_OUTPUT_COPY_TYPE = 4; /* type info: OutputCopy */
273-const uint32_t MSPROF_REPORT_MODEL_LOGIC_STREAM_TYPE = 7; /* type info: logic_stream_info */294+ const uint32_t MSPROF_REPORT_MODEL_LOGIC_STREAM_TYPE = 7; /* type info: logic_stream_info */
274-const uint32_t MSPROF_REPORT_MODEL_EXEOM_TYPE = 8; /* type info: exeom */295+ const uint32_t MSPROF_REPORT_MODEL_EXEOM_TYPE = 8; /* type info: exeom */
275-const uint32_t MSPROF_REPORT_MODEL_UDF_BASE_TYPE = 0x010000U; /* type info: udf_info */296+ const uint32_t MSPROF_REPORT_MODEL_UDF_BASE_TYPE = 0x010000U; /* type info: udf_info */
276-const uint32_t MSPROF_REPORT_MODEL_AICPU_BASE_TYPE = 0x020000U; /* type info: aicpu */297+ const uint32_t MSPROF_REPORT_MODEL_AICPU_BASE_TYPE = 0x020000U; /* type info: aicpu */
277 298 
278-/* Msprof report type of node(10000) level, offset: 0x000000 */299+ /* Msprof report type of node(10000) level, offset: 0x000000 */
279-const uint32_t MSPROF_REPORT_NODE_BASIC_INFO_TYPE = 0; /* type info: node_basic_info */300+ const uint32_t MSPROF_REPORT_NODE_BASIC_INFO_TYPE = 0; /* type info: node_basic_info */
280-const uint32_t MSPROF_REPORT_NODE_TENSOR_INFO_TYPE = 1; /* type info: tensor_info */301+ const uint32_t MSPROF_REPORT_NODE_TENSOR_INFO_TYPE = 1; /* type info: tensor_info */
281-const uint32_t MSPROF_REPORT_NODE_FUSION_OP_INFO_TYPE = 2; /* type info: fusion_op_info */302+ const uint32_t MSPROF_REPORT_NODE_FUSION_OP_INFO_TYPE = 2; /* type info: fusion_op_info */
282-const uint32_t MSPROF_REPORT_NODE_CONTEXT_ID_INFO_TYPE = 4; /* type info: context_id_info */303+ const uint32_t MSPROF_REPORT_NODE_CONTEXT_ID_INFO_TYPE = 4; /* type info: context_id_info */
283-const uint32_t MSPROF_REPORT_NODE_LAUNCH_TYPE = 5; /* type info: launch */304+ const uint32_t MSPROF_REPORT_NODE_LAUNCH_TYPE = 5; /* type info: launch */
284-const uint32_t MSPROF_REPORT_NODE_TASK_MEMORY_TYPE = 6; /* type info: task_memory_info */305+ const uint32_t MSPROF_REPORT_NODE_TASK_MEMORY_TYPE = 6; /* type info: task_memory_info */
285-const uint32_t MSPROF_REPORT_NODE_HOST_OP_EXEC_TYPE = 8; /* type info: op exec */306+ const uint32_t MSPROF_REPORT_NODE_HOST_OP_EXEC_TYPE = 8; /* type info: op exec */
286-const uint32_t MSPROF_REPORT_NODE_ATTR_INFO_TYPE = 9; /* type info: node_attr_info */307+ const uint32_t MSPROF_REPORT_NODE_ATTR_INFO_TYPE = 9; /* type info: node_attr_info */
287-const uint32_t MSPROF_REPORT_NODE_HCCL_OP_INFO_TYPE = 10; /* type info: hccl op info */308+ const uint32_t MSPROF_REPORT_NODE_HCCL_OP_INFO_TYPE = 10; /* type info: hccl op info */
288-const uint32_t MSPROF_REPORT_NODE_STATIC_OP_MEM_TYPE = 11; /* type info: static_op_mem */309+ const uint32_t MSPROF_REPORT_NODE_STATIC_OP_MEM_TYPE = 11; /* type info: static_op_mem */
289-/* Msprof report type of node(10000) level(ge api), offset: 0x010000 */310+ /* Msprof report type of node(10000) level(ge api), offset: 0x010000 */
290-const uint32_t MSPROF_REPORT_NODE_GE_API_BASE_TYPE = 0x010000U;311+ const uint32_t MSPROF_REPORT_NODE_GE_API_BASE_TYPE = 0x010000U;
291-const uint32_t MSPROF_REPORT_NODE_HCCL_BASE_TYPE = 0x020000U; /* type info: hccl api */312+ const uint32_t MSPROF_REPORT_NODE_HCCL_BASE_TYPE = 0x020000U; /* type info: hccl api */
292-const uint32_t MSPROF_REPORT_NODE_DVPP_API_BASE_TYPE = 0x030000U; /* type info: dvpp api */313+ const uint32_t MSPROF_REPORT_NODE_DVPP_API_BASE_TYPE = 0x030000U; /* type info: dvpp api */
293 314 
294-/* Msprof report type of hccl(5500) level(op api), offset: 0x010000 */315+ /* Msprof report type of hccl(5500) level(op api), offset: 0x010000 */
295-const uint32_t MSPROF_REPORT_HCCL_NODE_BASE_TYPE = 0x010000U;316+ const uint32_t MSPROF_REPORT_HCCL_NODE_BASE_TYPE = 0x010000U;
296-const uint32_t MSPROF_REPORT_HCCL_MASTER_TYPE = 0x010001U;317+ const uint32_t MSPROF_REPORT_HCCL_MASTER_TYPE = 0x010001U;
297-const uint32_t MSPROF_REPORT_HCCL_SLAVE_TYPE = 0x010002U;318+ const uint32_t MSPROF_REPORT_HCCL_SLAVE_TYPE = 0x010002U;
298 319 
299-struct MsprofApi { // for MsprofReportApi320+ struct MsprofApi
300- uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;321+ { // for MsprofReportApi
301- uint16_t level;322+ uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
302- uint32_t type;323+ uint16_t level;
303- uint32_t threadId;324+ uint32_t type;
304- uint32_t reserve;325+ uint32_t threadId;
305- uint64_t beginTime;326+ uint32_t reserve;
306- uint64_t endTime;327+ uint64_t beginTime;
307- uint64_t itemId;328+ uint64_t endTime;
308-};329+ uint64_t itemId;
330+ };
309 331 
310-struct MsprofEvent { // for MsprofReportEvent332+ struct MsprofEvent
311- uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;333+ { // for MsprofReportEvent
312- uint16_t level;334+ uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
313- uint32_t type;335+ uint16_t level;
314- uint32_t threadId;336+ uint32_t type;
315- uint32_t requestId; // 0xFFFF means single event337+ uint32_t threadId;
316- uint64_t timeStamp;338+ uint32_t requestId; // 0xFFFF means single event
317- uint64_t reserve = MSPROF_EVENT_FLAG;339+ uint64_t timeStamp;
318- uint64_t itemId;340+ uint64_t reserve = MSPROF_EVENT_FLAG;
319-};341+ uint64_t itemId;
342+ };
320 343 
321-struct MsprofRuntimeTrack { // for MsprofReportCompactInfo buffer data344+ struct MsporfKernelInfo
322- uint16_t deviceId;345+ {
323- uint16_t streamId;346+ uint16_t numBlocks;
324- uint32_t taskId;347+ uint16_t argsSize;
325- uint64_t taskType;348+ uint8_t ratio : 3;
326- uint64_t kernelName;349+ uint8_t schedMode : 2;
327-};350+ uint8_t rsv : 3;
351+ uint8_t reserved[11];
352+ };
328 353 
329-struct MsprofCaptureStreamInfo { // for MsprofReportCompactInfo buffer data354+ struct MsprofDim3
330- uint16_t captureStatus; // 标志是否销毁 0记为正常 1记为销毁355+ {
331- uint16_t modelStreamId; // capture stream id 销毁记录的stream id设置为 UINT16_MAX356+ uint16_t x;
332- uint16_t originalStreamId; // ori stream id 销毁记录的stream id设置为 UINT16_MAX357+ uint16_t y;
333- uint16_t modelId; // capture model id与GE无关358+ uint16_t z;
334- uint16_t deviceId;359+ };
335-};
336 360 
337-enum AlgType {361+ struct MsprofSimtKernelInfo
338- HCCL_ALG_NONE = 0,362+ {
339- HCCL_ALG_MESH,363+ MsprofDim3 gridDim;
340- HCCL_ALG_RING,364+ MsprofDim3 blockDim;
341- HCCL_ALG_NB,365+ uint16_t argsSize;
342- HCCL_ALG_HD,366+ uint8_t schedMode : 2;
343- HCCL_ALG_NHR,367+ uint8_t rsv : 6;
344- HCCL_ALG_PIPELINE,368+ uint8_t reserved;
345- HCCL_ALG_PAIRWISE,369+ };
346- HCCL_ALG_STAR,370+ 
347-};371+ struct MsprofRuntimeTrack
372+ { // for MsprofReportCompactInfo buffer data
373+ uint16_t deviceId;
374+ uint16_t streamId;
375+ uint32_t taskId;
376+ uint64_t taskType;
377+ uint64_t kernelName;
378+ union
379+ {
380+ struct MsporfKernelInfo kernelInfo;
381+ struct MsprofSimtKernelInfo simtKernelInfo;
382+ } extInfo;
383+ };
384+ 
385+ struct MsprofCaptureStreamInfo
386+ { // for MsprofReportCompactInfo buffer data
387+ uint16_t captureStatus; // 标志是否销毁 0记为正常 1记为销毁
388+ uint16_t modelStreamId; // capture stream id 销毁记录的stream id设置为 UINT16_MAX
389+ uint16_t originalStreamId; // ori stream id 销毁记录的stream id设置为 UINT16_MAX
390+ uint16_t modelId; // capture model id与GE无关
391+ uint16_t deviceId;
392+ };
393+ 
394+ enum AlgType
395+ {
396+ HCCL_ALG_NONE = 0,
397+ HCCL_ALG_MESH,
398+ HCCL_ALG_RING,
399+ HCCL_ALG_NB,
400+ HCCL_ALG_HD,
401+ HCCL_ALG_NHR,
402+ HCCL_ALG_PIPELINE,
403+ HCCL_ALG_PAIRWISE,
404+ HCCL_ALG_STAR,
405+ };
348 406 
349#pragma pack(1)407#pragma pack(1)
350-struct MsprofHcclOPInfo { // for MsprofReportCompactInfo buffer data408+ struct MsprofHcclOPInfo
351- uint8_t relay : 1;409+ { // for MsprofReportCompactInfo buffer data
352- uint8_t retry : 1;410+ uint8_t relay : 1;
353- uint8_t dataType;411+ uint8_t retry : 1;
354- uint64_t algType; // 通信算子使用的算法,hash的key,其值是以"-"分隔的字符串412+ uint8_t dataType;
355- uint64_t count;413+ uint64_t algType; // 通信算子使用的算法,hash的key,其值是以"-"分隔的字符串
356- uint64_t groupName;414+ uint64_t count;
357-};415+ uint64_t groupName;
416+ };
358#pragma pack()417#pragma pack()
359 418 
360-struct MsprofMemcpyInfo { // for MsprofReportCompactInfo buffer data419+ struct MsprofMemcpyInfo
361- uint64_t dataSize; // 数据量大小， 字节420+ { // for MsprofReportCompactInfo buffer data
362- uint64_t maxSize; // 单个task拷贝的最大数据量421+ uint64_t dataSize; // 数据量大小， 字节
363- uint16_t memcpyDirection; // memcpy的方向422+ uint64_t maxSize; // 单个task拷贝的最大数据量
364-};423+ uint16_t memcpyDirection; // memcpy的方向
365- 
366-const uint16_t MSPROF_AICPU_DATA_RESERVE_BYTES = 9;
367-struct MsprofAicpuNodeAdditionalData {
368- uint16_t streamId;
369- uint16_t taskId;
370- uint32_t rev;
371- uint64_t runStartTime;
372- uint64_t runStartTick;
373- uint64_t computeStartTime;
374- uint64_t memcpyStartTime;
375- uint64_t memcpyEndTime;
376- uint64_t runEndTime;
377- uint64_t runEndTick;
378- uint32_t threadId;
379- uint32_t deviceId;
380- uint64_t submitTick;
381- uint64_t scheduleTick;
382- uint64_t tickBeforeRun;
383- uint64_t tickAfterRun;
384- uint32_t kernelType;
385- uint32_t dispatchTime;
386- uint32_t totalTime;
387- uint16_t fftsThreadId;
388- uint8_t version;
389- uint8_t reserve[MSPROF_AICPU_DATA_RESERVE_BYTES];
390-};
391- 
392-const uint16_t MSPROF_AICPU_MODEL_RESERVE_BYTES = 24;
393-struct MsprofAicpuModelAdditionalData {
394- uint64_t indexId;
395- uint32_t modelId;
396- uint16_t tagId;
397- uint16_t rsv1;
398- uint64_t eventId;
399- uint8_t reserve[MSPROF_AICPU_MODEL_RESERVE_BYTES];
400-};
401- 
402-const uint16_t MSPROF_DP_DATA_RESERVE_BYTES = 16;
403-const uint16_t MSPROF_DP_DATA_ACTION_LEN = 16;
404-const uint16_t MSPROF_DP_DATA_SOURCE_LEN = 64;
405-struct MsprofAicpuDpAdditionalData {
406- char action[MSPROF_DP_DATA_ACTION_LEN];
407- char source[MSPROF_DP_DATA_SOURCE_LEN];
408- uint64_t index;
409- uint64_t size;
410- uint8_t reserve[MSPROF_DP_DATA_RESERVE_BYTES];
411-};
412- 
413-enum MsprofMindsporeNodeTag {
414- GET_NEXT_DEQUEUE_WAIT = 1,
415-};
416- 
417-struct MsprofAicpuMiAdditionalData {
418- uint32_t nodeTag; // MsprofMindsporeNodeTag:1
419- uint32_t reserve;
420- uint64_t queueSize;
421- uint64_t runStartTime;
422- uint64_t runEndTime;
423-};
424- 
425-const uint16_t MSPROF_COMPACT_INFO_DATA_LENGTH = 40;
426-struct MsprofCompactInfo { // for MsprofReportCompactInfo buffer data
427- uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
428- uint16_t level;
429- uint32_t type;
430- uint32_t threadId;
431- uint32_t dataLen;
432- uint64_t timeStamp;
433- union {
434- uint8_t info[MSPROF_COMPACT_INFO_DATA_LENGTH];
435- MsprofRuntimeTrack runtimeTrack;
436- MsprofCaptureStreamInfo captureStreamInfo;
437- MsprofNodeBasicInfo nodeBasicInfo;
438- MsprofAttrInfo nodeAttrInfo;
439- MsprofHcclOPInfo hcclopInfo;
440- MsprofMemcpyInfo memcpyInfo;
441- } data;
442-};
443- 
444-const uint16_t MSPROF_ADDITIONAL_INFO_DATA_LENGTH = 232;
445-struct MsprofAdditionalInfo { // for MsprofReportAdditionalInfo buffer data
446- uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
447- uint16_t level;
448- uint32_t type;
449- uint32_t threadId;
450- uint32_t dataLen;
451- uint64_t timeStamp;
452- union {
453- uint8_t data[MSPROF_ADDITIONAL_INFO_DATA_LENGTH];
454- MsprofAicpuNodeAdditionalData aicpuNode;
455- MsprofAicpuModelAdditionalData aicpuModel;
456- MsprofAicpuDpAdditionalData aicpuDp;
457- MsprofAicpuMiAdditionalData aicpuMi;
458 };424 };
459-};
460 425 
461-struct MsprofVariableInfo { // for MsprofVariableInfo buffer data426+ const uint16_t MSPROF_AICPU_DATA_RESERVE_BYTES = 9;
462- uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;427+ struct MsprofAicpuNodeAdditionalData
463- uint16_t level;428+ {
464- uint32_t type;429+ uint16_t streamId;
465- uint32_t threadId;430+ uint16_t taskId;
466- uint32_t dataLen;431+ uint32_t rev;
467- uint64_t timeStamp;432+ uint64_t runStartTime;
468- uint8_t data[0];433+ uint64_t runStartTick;
469-};434+ uint64_t computeStartTime;
435+ uint64_t memcpyStartTime;
436+ uint64_t memcpyEndTime;
437+ uint64_t runEndTime;
438+ uint64_t runEndTick;
439+ uint32_t threadId;
440+ uint32_t deviceId;
441+ uint64_t submitTick;
442+ uint64_t scheduleTick;
443+ uint64_t tickBeforeRun;
444+ uint64_t tickAfterRun;
445+ uint32_t kernelType;
446+ uint32_t dispatchTime;
447+ uint32_t totalTime;
448+ uint16_t fftsThreadId;
449+ uint8_t version;
450+ uint8_t reserve[MSPROF_AICPU_DATA_RESERVE_BYTES];
451+ };
470 452 
471-struct ConcatTensorInfo {453+ const uint16_t MSPROF_AICPU_MODEL_RESERVE_BYTES = 24;
472- uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;454+ struct MsprofAicpuModelAdditionalData
473- uint16_t level = 0;455+ {
474- uint32_t type = 0;456+ uint64_t indexId;
475- uint32_t threadId = 0;457+ uint32_t modelId;
476- uint32_t dataLen = 0;458+ uint16_t tagId;
477- uint64_t timeStamp = 0;459+ uint16_t rsv1;
478- uint64_t opName = 0;460+ uint64_t eventId;
479- uint32_t tensorNum = 0;461+ uint8_t reserve[MSPROF_AICPU_MODEL_RESERVE_BYTES];
480- std::vector<MsrofTensorData> tensorData{MSPROF_GE_TENSOR_DATA_NUM};462+ };
481-};
482 463 
483-// AICPU kfc算子执行时间464+ const uint16_t MSPROF_DP_DATA_RESERVE_BYTES = 16;
484-struct AicpuKfcProfCommTurn {465+ const uint16_t MSPROF_DP_DATA_ACTION_LEN = 16;
485- uint64_t serverStartTime; // 进入KFC流程466+ const uint16_t MSPROF_DP_DATA_SOURCE_LEN = 64;
486- uint64_t waitMsgStartTime; // 开始等待客户端消息467+ struct MsprofAicpuDpAdditionalData
487- uint64_t kfcAlgExeStartTime; // 开始通信算法执行468+ {
488- uint64_t sendTaskStartTime; // 开始下发task469+ char action[MSPROF_DP_DATA_ACTION_LEN];
489- uint64_t sendSqeFinishTime; // task下发完成470+ char source[MSPROF_DP_DATA_SOURCE_LEN];
490- uint64_t rtsqExeEndTime; // sq执行结束时间471+ uint64_t index;
491- uint64_t serverEndTime; // KFC流程结束时间472+ uint64_t size;
492- uint64_t dataLen; // 本轮通信数据长度473+ uint8_t reserve[MSPROF_DP_DATA_RESERVE_BYTES];
493- uint32_t deviceId;474+ };
494- uint16_t streamId;475+ 
495- uint16_t taskId;476+ enum MsprofMindsporeNodeTag
496- uint8_t version;477+ {
497- uint8_t commTurn; // 总通信轮次478+ GET_NEXT_DEQUEUE_WAIT = 1,
498- uint8_t currentTurn;479+ };
499- uint8_t reserve[5];480+ 
500-};481+ struct MsprofAicpuMiAdditionalData
482+ {
483+ uint32_t nodeTag; // MsprofMindsporeNodeTag:1
484+ uint32_t reserve;
485+ uint64_t queueSize;
486+ uint64_t runStartTime;
487+ uint64_t runEndTime;
488+ };
489+ 
490+ const uint16_t MSPROF_COMPACT_INFO_DATA_LENGTH = 40;
491+ struct MsprofCompactInfo
492+ { // for MsprofReportCompactInfo buffer data
493+ uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
494+ uint16_t level;
495+ uint32_t type;
496+ uint32_t threadId;
497+ uint32_t dataLen;
498+ uint64_t timeStamp;
499+ union
500+ {
501+ uint8_t info[MSPROF_COMPACT_INFO_DATA_LENGTH];
502+ MsprofRuntimeTrack runtimeTrack;
503+ MsprofCaptureStreamInfo captureStreamInfo;
504+ MsprofNodeBasicInfo nodeBasicInfo;
505+ MsprofAttrInfo nodeAttrInfo;
506+ MsprofHcclOPInfo hcclopInfo;
507+ MsprofMemcpyInfo memcpyInfo;
508+ } data;
509+ };
510+ 
511+ const uint16_t MSPROF_ADDITIONAL_INFO_DATA_LENGTH = 232;
512+ struct MsprofAdditionalInfo
513+ { // for MsprofReportAdditionalInfo buffer data
514+ uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
515+ uint16_t level;
516+ uint32_t type;
517+ uint32_t threadId;
518+ uint32_t dataLen;
519+ uint64_t timeStamp;
520+ union
521+ {
522+ uint8_t data[MSPROF_ADDITIONAL_INFO_DATA_LENGTH];
523+ MsprofAicpuNodeAdditionalData aicpuNode;
524+ MsprofAicpuModelAdditionalData aicpuModel;
525+ MsprofAicpuDpAdditionalData aicpuDp;
526+ MsprofAicpuMiAdditionalData aicpuMi;
527+ };
528+ };
529+ 
530+ struct MsprofVariableInfo
531+ { // for MsprofVariableInfo buffer data
532+ uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
533+ uint16_t level;
534+ uint32_t type;
535+ uint32_t threadId;
536+ uint32_t dataLen;
537+ uint64_t timeStamp;
538+ uint8_t data[0];
539+ };
540+ 
541+ struct ConcatTensorInfo
542+ {
543+ uint16_t magicNumber = MSPROF_DATA_HEAD_MAGIC_NUM;
544+ uint16_t level = 0;
545+ uint32_t type = 0;
546+ uint32_t threadId = 0;
547+ uint32_t dataLen = 0;
548+ uint64_t timeStamp = 0;
549+ uint64_t opName = 0;
550+ uint32_t tensorNum = 0;
551+ std::vector<MsrofTensorData> tensorData{MSPROF_GE_TENSOR_DATA_NUM};
552+ };
553+ 
554+ // AICPU kfc算子执行时间
555+ struct AicpuKfcProfCommTurn
556+ {
557+ uint64_t serverStartTime; // 进入KFC流程
558+ uint64_t waitMsgStartTime; // 开始等待客户端消息
559+ uint64_t kfcAlgExeStartTime; // 开始通信算法执行
560+ uint64_t sendTaskStartTime; // 开始下发task
561+ uint64_t sendSqeFinishTime; // task下发完成
562+ uint64_t rtsqExeEndTime; // sq执行结束时间
563+ uint64_t serverEndTime; // KFC流程结束时间
564+ uint64_t dataLen; // 本轮通信数据长度
565+ uint32_t deviceId;
566+ uint16_t streamId;
567+ uint16_t taskId;
568+ uint8_t version;
569+ uint8_t commTurn; // 总通信轮次
570+ uint8_t currentTurn;
571+ uint8_t reserve[5];
572+ };
501 573 
502#ifdef __cplusplus574#ifdef __cplusplus
503}575}
Manalysis/csrc/viewer/database/finals/msprof_db.cpp+3-1
@@ -60,7 +60,9 @@ const TableColumns COMPUTE_TASK_INFO = {{"name", SQL_INTEGER_TYPE},
60 {"outputShapes", SQL_INTEGER_TYPE},60 {"outputShapes", SQL_INTEGER_TYPE},
61 {"attrInfo", SQL_INTEGER_TYPE},61 {"attrInfo", SQL_INTEGER_TYPE},
62 {"opState", SQL_INTEGER_TYPE},62 {"opState", SQL_INTEGER_TYPE},
63- {"hf32Eligible", SQL_INTEGER_TYPE}};63+ {"hf32Eligible", SQL_INTEGER_TYPE},
64+ {"gridDim", SQL_INTEGER_TYPE},
65+ {"blockDim", SQL_INTEGER_TYPE}};
W
Wwangzixuan6月1日

[review] 这里新增的字段也叫blockDim。但是之前表字段中刚从blockDim变更为blockNum。这里新加的字段和以前的重复，存在版本上的含义变更和不兼容。不建议重新设置为blockDim。可考虑其他命名

likedislike
xfeng
xfeng
6月1日 评论：
64 66 
65const TableColumns COMMUNICATION_SCHEDULE_TASK_INFO = {{"name", SQL_INTEGER_TYPE},67const TableColumns COMMUNICATION_SCHEDULE_TASK_INFO = {{"name", SQL_INTEGER_TYPE},
66 {"globalTaskId", SQL_INTEGER_TYPE, true},68 {"globalTaskId", SQL_INTEGER_TYPE, true},
Manalysis/msmodel/ge/ge_info_model.py+34-10
@@ -16,7 +16,6 @@
16from collections import namedtuple16from collections import namedtuple
17 17 
18from common_func.db_name_constant import DBNameConstant18from common_func.db_name_constant import DBNameConstant
19-from common_func.constant import Constant
20from common_func.db_manager import DBManager19from common_func.db_manager import DBManager
21from common_func.msprof_object import CustomizedNamedtupleFactory20from common_func.msprof_object import CustomizedNamedtupleFactory
22from msmodel.interface.parser_model import ParserModel21from msmodel.interface.parser_model import ParserModel
@@ -52,7 +51,7 @@ class GeModel(ParserModel):
52 """51 """
53 delete ge data52 delete ge data
54 """53 """
55- self.cur.execute('delete from {}'.format(table_name))54+ self.cur.execute('delete from {}'.format(table_name)) # nosec
56 55 
57 def get_ge_model_name(self: any) -> any:56 def get_ge_model_name(self: any) -> any:
58 """57 """
@@ -63,19 +62,44 @@ class GeModel(ParserModel):
63 62 
64class GeInfoViewModel(ViewModel):63class GeInfoViewModel(ViewModel):
65 TASK_INFO_TYPE = CustomizedNamedtupleFactory.enhance_namedtuple(64 TASK_INFO_TYPE = CustomizedNamedtupleFactory.enhance_namedtuple(
66- namedtuple("TaskInfo",65+ namedtuple(
67- ["model_id", "op_name", "stream_id", "task_id", "block_num", "mix_block_num",66+ "TaskInfo",
68- "op_stat", "task_type", "op_type", "index_id", "thread_id", "timestamp", "batch_id",67+ [
69- "tensor_num", "input_formats", "input_data_types", "input_shapes",68+ "model_id",
70- "output_formats", "output_data_types", "output_shapes", "device_id", "context_id",69+ "op_name",
71- "op_flag", "hashid"]),70+ "stream_id",
72- {})71+ "task_id",
72+ "block_num",
73+ "mix_block_num",
74+ "op_state",
75+ "task_type",
76+ "op_type",
77+ "index_id",
78+ "thread_id",
79+ "timestamp",
80+ "batch_id",
81+ "tensor_num",
82+ "input_formats",
83+ "input_data_types",
84+ "input_shapes",
85+ "output_formats",
86+ "output_data_types",
87+ "output_shapes",
88+ "device_id",
89+ "context_id",
90+ "op_flag",
91+ "hashid",
92+ ],
93+ ),
94+ {},
95+ )
73 96 
74 def __init__(self, result_dir: str, table_list: list):97 def __init__(self, result_dir: str, table_list: list):
75 super().__init__(result_dir, DBNameConstant.DB_GE_INFO, table_list)98 super().__init__(result_dir, DBNameConstant.DB_GE_INFO, table_list)
76 99 
77 def get_ge_info_by_device_id(self: any, table_name: str, device_id: str, task_type_filter: tuple = tuple()) -> any:100 def get_ge_info_by_device_id(self: any, table_name: str, device_id: str, task_type_filter: tuple = tuple()) -> any:
78- ge_sql = "select * from {0} where device_id={1} ".format(table_name, device_id)101+ fields = ", ".join(self.TASK_INFO_TYPE._fields)
102+ ge_sql = "select {0} from {1} where device_id={2} ".format(fields, table_name, device_id) # nosec
79 condition = ""103 condition = ""
80 for t in task_type_filter:104 for t in task_type_filter:
81 condition += " AND task_type != '{0}' ".format(t)105 condition += " AND task_type != '{0}' ".format(t)
Mtest/msprof_cpp/analysis_ut/application/assemble/test/ascend_hardware_assembler_utest.cpp+70-0
@@ -180,6 +180,43 @@ static std::vector<MemcpyInfoData> GenerateMemcpyInfoData()
180 return res;180 return res;
181}181}
182 182 
183+static std::vector<AscendTaskData> GenerateSimtTaskData()
184+{
185+ std::vector<AscendTaskData> res;
186+ AscendTaskData data;
187+ data.deviceId = 0; // deviceId 0
188+ data.indexId = -1; // index_id -1
189+ data.streamId = 1; // streamId 1
190+ data.taskId = 10; // taskId 10
191+ data.contextId = 1; // contextId 1
192+ data.batchId = 1; // batchId 1
193+ data.connectionId = 2345; // connectionId 2345
194+ data.timestamp = 1717575960208020758; // start 1717575960208020758
195+ data.duration = 450.78; // dur 450.78
196+ data.hostType = "KERNEL_SIMT";
197+ data.deviceType = "AI_CORE";
198+ data.taskType = "KERNEL_SIMT";
199+ res.push_back(data);
200+ return res;
201+}
202+ 
203+static std::vector<TaskInfoData> GenerateSimtTaskInfoData()
204+{
205+ std::vector<TaskInfoData> res;
206+ TaskInfoData data;
207+ data.deviceId = 0; // deviceId 0
208+ data.streamId = 1; // streamId 1
209+ data.taskId = 10; // taskId 10
210+ data.contextId = 1; // contextId 1
211+ data.batchId = 1; // batchId 1
212+ data.opName = "MatMulV3";
213+ data.taskType = "KERNEL_SIMT";
214+ data.gridDim = "2,3,4";
215+ data.blockDim = "5,6,7";
216+ res.push_back(data);
217+ return res;
218+}
219+ 
183TEST_F(AscendHardwareAssemblerUTest, ShouldReturnTrueWhenDataNotExists)220TEST_F(AscendHardwareAssemblerUTest, ShouldReturnTrueWhenDataNotExists)
184{221{
185 AscendHardwareAssembler assembler;222 AscendHardwareAssembler assembler;
@@ -430,4 +467,37 @@ TEST_F(AscendHardwareAssemblerUTest, ShouldReturnTrueWhenDataAssembleWithLogicSt
430 ":\"thread_sort_index\",\"pid\":10328512,\"tid\":1337,\"ph\":\"M\",\"args\""467 ":\"thread_sort_index\",\"pid\":10328512,\"tid\":1337,\"ph\":\"M\",\"args\""
431 ":{\"sort_index\":1337}},";468 ":{\"sort_index\":1337}},";
432 EXPECT_EQ(expStr, res.back());469 EXPECT_EQ(expStr, res.back());
470+}
471+ 
472+TEST_F(AscendHardwareAssemblerUTest, ShouldReturnTrueWhenDataAssembleWithSimtTaskType)
473+{
474+ AscendHardwareAssembler assembler;
475+ std::shared_ptr<std::vector<AscendTaskData>> taskS;
476+ std::shared_ptr<std::vector<TaskInfoData>> infoS;
477+ auto task = GenerateSimtTaskData();
478+ auto info = GenerateSimtTaskInfoData();
479+ MAKE_SHARED_NO_OPERATION(taskS, std::vector<AscendTaskData>, task);
480+ MAKE_SHARED_NO_OPERATION(infoS, std::vector<TaskInfoData>, info);
481+ dataInventory_.Inject(taskS);
482+ dataInventory_.Inject(infoS);
483+ MOCKER_CPP(&Context::GetPidFromInfoJson).stubs().will(returnValue(10086)); // pid 10086
484+ EXPECT_TRUE(assembler.Run(dataInventory_, PROF_PATH));
485+ auto files = File::GetOriginData(RESULT_PATH, {"msprof"}, {});
486+ EXPECT_EQ(1ul, files.size());
487+ FileReader reader(files.back());
488+ std::vector<std::string> res;
489+ EXPECT_EQ(Analysis::ANALYSIS_OK, reader.ReadText(res));
490+ std::string expectStr = "{\"name\":\"MatMulV3\",\"pid\":10328512,\"tid\":1,\"ts\":\"1717575960208020.758\",\"dur\""
491+ ":0.45077999999999996,\"ph\":\"X\",\"args\":{\"Model Id\":4294967295,\"Task Type\":\""
492+ "KERNEL_SIMT\",\"Physic Stream Id\":1,\"Task Id\":10,\"Batch Id\":1,\"Subtask Id\":1,\""
493+ "connection_id\":2345,\"Grid Dim\":\"2,3,4\",\"Block Dim\":\"5,6,7\"}},{\"name\":\""
494+ "HostToDevice10071698309120\",\"pid\":10328512,\"tid\":1,\"ph\":\"f\",\"cat\":\""
495+ "HostToDevice\",\"id\":\"10071698309120\",\"ts\":\"1717575960208020.758\",\"bp\":\""
496+ "e\"},{\"name\":\"process_name\",\"pid\":10328512,\"tid\":0,\"ph\":\"M\",\"args\":{\"name\""
497+ ":\"Ascend Hardware\"}},{\"name\":\"process_labels\",\"pid\":10328512,\"tid\":0,\"ph\":\"M\""
498+ ",\"args\":{\"labels\":\"NPU 0\"}},{\"name\":\"process_sort_index\",\"pid\":10328512,\"tid\""
499+ ":0,\"ph\":\"M\",\"args\":{\"sort_index\":14}},{\"name\":\"thread_name\",\"pid\":10328512,\""
500+ "tid\":1,\"ph\":\"M\",\"args\":{\"name\":\"Stream 1\"}},{\"name\":\"thread_sort_index\",\""
501+ "pid\":10328512,\"tid\":1,\"ph\":\"M\",\"args\":{\"sort_index\":1}},";
502+ EXPECT_EQ(expectStr, res.back());
433}503}
Mtest/msprof_cpp/analysis_ut/application/assemble/test/db_assembler_utest.cpp+1-1
@@ -1555,7 +1555,7 @@ TEST_F(DBAssemblerUTest, TestSaveComputeTaskInfoShouldReturnFalseWhenReserveFail
1555 using ComputeTaskInfoFormat = std::vector<std::tuple<uint64_t, uint64_t, uint32_t, uint32_t,1555 using ComputeTaskInfoFormat = std::vector<std::tuple<uint64_t, uint64_t, uint32_t, uint32_t,
1556 uint64_t, uint64_t, uint64_t, uint64_t,1556 uint64_t, uint64_t, uint64_t, uint64_t,
1557 uint64_t, uint64_t, uint64_t, uint64_t,1557 uint64_t, uint64_t, uint64_t, uint64_t,
1558- uint64_t, uint64_t, uint64_t>>;1558+ uint64_t, uint64_t, uint64_t, uint64_t, uint64_t>>;
1559 std::vector<TaskInfoData> res;1559 std::vector<TaskInfoData> res;
1560 TaskInfoData data;1560 TaskInfoData data;
1561 res.push_back(data);1561 res.push_back(data);
Mtest/msprof_cpp/analysis_ut/domain/data_process/test/compute_task_info_processor_utest.cpp+3-3
@@ -42,14 +42,14 @@ const std::string TABLE_NAME = "TaskInfo";
42using GeInfoFormat = std::vector<std::tuple<uint32_t, std::string, int32_t, int32_t, uint32_t, uint32_t, std::string,42using GeInfoFormat = std::vector<std::tuple<uint32_t, std::string, int32_t, int32_t, uint32_t, uint32_t, std::string,
43 std::string, std::string, int32_t, uint32_t, double, uint32_t, uint32_t,43 std::string, std::string, int32_t, uint32_t, double, uint32_t, uint32_t,
44 std::string, std::string, std::string, std::string, std::string,44 std::string, std::string, std::string, std::string, std::string,
45- std::string, int32_t, uint32_t, std::string, std::string>>;45+ std::string, int32_t, uint32_t, std::string, std::string, std::string, std::string>>;
46 46 
47GeInfoFormat DATA_A{{4294967295, "aclnnMm_MatMulCommon_MatMulV2", 2, 1, 20, 40, "1", "MIX_AIC", "MatMulV2", -1,47GeInfoFormat DATA_A{{4294967295, "aclnnMm_MatMulCommon_MatMulV2", 2, 1, 20, 40, "1", "MIX_AIC", "MatMulV2", -1,
48 3391981, 453148218443103, 0, 3, "FORMAT_ND;FORMAT_ND", "FLOAT16;FLOAT16",48 3391981, 453148218443103, 0, 3, "FORMAT_ND;FORMAT_ND", "FLOAT16;FLOAT16",
49- "\"10000,10000;10000,10000\"", "FORMAT_ND", "FLOAT16", "\"1000,1000\"", 0, 0, "NO", "N/A"},49+ "\"10000,10000;10000,10000\"", "FORMAT_ND", "FLOAT16", "\"1000,1000\"", 0, 0, "NO", "N/A", "N/A", "N/A"},
50 {4294967295, "trans_TransData_0", 2, 2, 35, 0, "1", "AI_CORE", "TransData", -1,50 {4294967295, "trans_TransData_0", 2, 2, 35, 0, "1", "AI_CORE", "TransData", -1,
51 250512, 569402956566, 0, 2, "FORMAT_ND", "FLOAT",51 250512, 569402956566, 0, 2, "FORMAT_ND", "FLOAT",
52- "\"3072,768\"", "FRACTAL_NZ", "FLOAT", "\"48,192,16,16\"", 0, 4294967295, "NO", "N/A"}};52+ "\"3072,768\"", "FRACTAL_NZ", "FLOAT", "\"48,192,16,16\"", 0, 4294967295, "NO", "N/A", "N/A", "N/A"}};
53}53}
54 54 
55class ComputeTaskInfoProcessorUTest : public testing::Test {55class ComputeTaskInfoProcessorUTest : public testing::Test {
Mtest/msprof_cpp/analysis_ut/domain/services/test/load_host_data_utest.cpp+4-2
@@ -78,6 +78,8 @@ using OriDataFormat = std::vector<
78 uint32_t,78 uint32_t,
79 uint32_t,79 uint32_t,
80 std::string,80 std::string,
81+ std::string,
82+ std::string,
81 std::string>>;83 std::string>>;
82using HcclOpOriDataFormat = std::vector<84using HcclOpOriDataFormat = std::vector<
83 std::tuple<85 std::tuple<
@@ -136,10 +138,10 @@ const std::string HCCLDB_PATH = File::PathJoin({HOST_PATH, "sqlite", "hccl.db"})
136const OriDataFormat DATA_A{138const OriDataFormat DATA_A{
137 {0, "Default/network/network/bert/bert/bert_embedding_postprocessor/StridedSliceD-op6673",139 {0, "Default/network/network/bert/bert/bert_embedding_postprocessor/StridedSliceD-op6673",
138 1, 1, 32, 0, 0, "AI_CORE", "StridedSliceD", 0, 120040, 458597374830, 1, 2, "DEFAULT_",140 1, 1, 32, 0, 0, "AI_CORE", "StridedSliceD", 0, 120040, 458597374830, 1, 2, "DEFAULT_",
139- "INT32_", "1,512", "DEFAULT_", "INT32_", "1,512", 0, 1, "1", "123"},141+ "INT32_", "1,512", "DEFAULT_", "INT32_", "1,512", 0, 1, "1", "123", "N/A", "N/A"},
140 {0, "Default/network/network/bert/bert/bert_embedding_postprocessor/StridedSliceD-op6673",142 {0, "Default/network/network/bert/bert/bert_embedding_postprocessor/StridedSliceD-op6673",
141 1, 2, 32, 0, 0, "AI_CORE", "StridedSliceD", 0, 120040, 458597374830, 1, 2, "DEFAULT_",143 1, 2, 32, 0, 0, "AI_CORE", "StridedSliceD", 0, 120040, 458597374830, 1, 2, "DEFAULT_",
142- "INT32_", "1,512", "DEFAULT_", "INT32_", "1,512", 0, 1, "1", "123"}};144+ "INT32_", "1,512", "DEFAULT_", "INT32_", "1,512", 0, 1, "1", "123", "N/A", "N/A"}};
143const runTimeData DATA_B{{4294967295, -1, 65535, 0, 4294967295, 0, "PROFILING_ENABLE", "aclnn", 0, 397936887714, -1, 1},145const runTimeData DATA_B{{4294967295, -1, 65535, 0, 4294967295, 0, "PROFILING_ENABLE", "aclnn", 0, 397936887714, -1, 1},
144 {4294967295, -1, 65535, 0, 4294967295, 0, "PROFILING_ENABLE", "aclnn", 0, 397936887714, -1, 1}};146 {4294967295, -1, 65535, 0, 4294967295, 0, "PROFILING_ENABLE", "aclnn", 0, 397936887714, -1, 1}};
145const HcclOpOriDataFormat DATA_HCCL_OP{{0, 3, 0, 121639, "Default/network/AllReduce-op0", "HCCL", "HcomAllReduce",147const HcclOpOriDataFormat DATA_HCCL_OP{{0, 3, 0, 121639, "Default/network/AllReduce-op0", "HCCL", "HcomAllReduce",
Mtest/msprof_cpp/analysis_ut/viewer/test/cann_db_dumper_utest.cpp+84-26
@@ -22,13 +22,16 @@
22 22 
23#include "analysis/csrc/infrastructure/utils/thread_pool.h"23#include "analysis/csrc/infrastructure/utils/thread_pool.h"
24#include "analysis/csrc/domain/services/persistence/host/cann_trace_db_dumper.h"24#include "analysis/csrc/domain/services/persistence/host/cann_trace_db_dumper.h"
25+#include "analysis/csrc/domain/services/environment/context.h"
25 26 
26 27 
27using namespace Analysis::Utils;28using namespace Analysis::Utils;
28using namespace Analysis::Infra;29using namespace Analysis::Infra;
29using namespace Analysis::Viewer::Database;30using namespace Analysis::Viewer::Database;
30using namespace Analysis::Domain;31using namespace Analysis::Domain;
32+using namespace Analysis::Domain::Environment;
31using namespace Analysis::Domain::Cann;33using namespace Analysis::Domain::Cann;
34+using TypeData = Analysis::Domain::Host::Cann::TypeData;
32const std::string TEST_DB_FILE_PATH = "./sqlite";35const std::string TEST_DB_FILE_PATH = "./sqlite";
33const uint32_t INPUT_DATA_TYPE_POSITION = 15;36const uint32_t INPUT_DATA_TYPE_POSITION = 15;
34const uint32_t GROUP_NAME_POSITION = 3;37const uint32_t GROUP_NAME_POSITION = 3;
@@ -107,6 +110,37 @@ protected:
107 MOCKER_CPP(&TreeAnalyzer::GetComputeTasks).stubs().will(returnValue(*kernelTasks));110 MOCKER_CPP(&TreeAnalyzer::GetComputeTasks).stubs().will(returnValue(*kernelTasks));
108 }111 }
109 112 
113+ static void MockGetComputeTasksWithMixedRuntimeTracks()
114+ {
115+ auto kernelTasks = std::make_shared<HostTasks>();
116+ auto makeTask = [&](uint64_t taskType, bool isSimt, uint16_t numBlocks, uint8_t ratio,
117+ uint16_t gx, uint16_t gy, uint16_t gz, uint16_t bx, uint16_t by, uint16_t bz) {
118+ auto kernelTask = std::make_shared<HostTask>();
119+ auto kernelDesc = std::make_shared<OpDesc>();
120+ kernelDesc->tensorDesc = std::make_shared<ConcatTensorInfo>();
121+ kernelDesc->nodeDesc = std::make_shared<MsprofCompactInfo>();
122+ kernelDesc->nodeDesc->data.nodeBasicInfo = MsprofNodeBasicInfo{};
123+ kernelDesc->nodeDesc->data.nodeAttrInfo = MsprofAttrInfo{};
124+ auto runtimeTrackDesc = std::make_shared<MsprofCompactInfo>();
125+ runtimeTrackDesc->dataLen = MSPROF_COMPACT_INFO_DATA_LENGTH;
126+ runtimeTrackDesc->data.runtimeTrack.taskType = taskType;
127+ if (isSimt) {
128+ runtimeTrackDesc->data.runtimeTrack.extInfo.simtKernelInfo.gridDim = {gx, gy, gz};
129+ runtimeTrackDesc->data.runtimeTrack.extInfo.simtKernelInfo.blockDim = {bx, by, bz};
130+ } else {
131+ runtimeTrackDesc->data.runtimeTrack.extInfo.kernelInfo.numBlocks = numBlocks;
132+ runtimeTrackDesc->data.runtimeTrack.extInfo.kernelInfo.ratio = ratio;
133+ }
134+ kernelDesc->runtimeTrackDesc = runtimeTrackDesc;
135+ kernelDesc->ctxId = std::make_shared<MsprofAdditionalInfo>();
136+ kernelTask->op = std::make_shared<Operator>(kernelDesc, 0, OpType::OPTYPE_COMPUTE);
137+ kernelTasks->push_back(std::make_shared<HostTask>(*kernelTask));
138+ };
139+ makeTask(99, true, 0, 0, 2, 3, 4, 5, 6, 7);
140+ makeTask(0, false, 100, 3, 0, 0, 0, 0, 0, 0);
141+ MOCKER_CPP(&TreeAnalyzer::GetComputeTasks).stubs().will(returnValue(*kernelTasks));
142+ }
143+ 
110 static void MockGetTasks()144 static void MockGetTasks()
111 {145 {
112 auto pMiniOpDesc = std::make_shared<HcclSmallOpDesc>(0, 0, nullptr);146 auto pMiniOpDesc = std::make_shared<HcclSmallOpDesc>(0, 0, nullptr);
@@ -176,13 +210,12 @@ TEST_F(CannDBDumperUtest,
176 GEInfoDB geInfoDB;210 GEInfoDB geInfoDB;
177 std::string opDescDBPath = Utils::File::PathJoin({TEST_DB_FILE_PATH, geInfoDB.GetDBName()});211 std::string opDescDBPath = Utils::File::PathJoin({TEST_DB_FILE_PATH, geInfoDB.GetDBName()});
178 DBRunner opDescDBRunner(opDescDBPath);212 DBRunner opDescDBRunner(opDescDBPath);
179- std::vector<std::tuple<uint32_t, std::string, uint32_t, uint32_t, uint32_t, uint32_t, std::string, uint32_t,213+ CANNTraceDBDumper::TaskInfoData taskInfoData;
180- std::string, uint32_t, uint32_t, double, uint32_t, uint32_t, std::string,214+ opDescDBRunner.QueryData("select * from TaskInfo", taskInfoData);
181- std::string, std::string, std::string, std::string, std::string, uint32_t,215+ EXPECT_EQ(taskInfoData.size(), 1);
182- uint32_t, uint32_t>> TaskInfoData;216+ EXPECT_EQ(std::get<INPUT_DATA_TYPE_POSITION>(taskInfoData[0]), NA);
183- opDescDBRunner.QueryData("select * from TaskInfo", TaskInfoData);217+ EXPECT_EQ(std::get<24>(taskInfoData[0]), NA);
184- EXPECT_EQ(TaskInfoData.size(), 1);218+ EXPECT_EQ(std::get<25>(taskInfoData[0]), NA);
185- EXPECT_EQ(std::get<INPUT_DATA_TYPE_POSITION>(TaskInfoData[0]), "N/A");
186 219 
187 std::vector<std::tuple<uint32_t, uint32_t, std::string, std::string, uint32_t, std::string,220 std::vector<std::tuple<uint32_t, uint32_t, std::string, std::string, uint32_t, std::string,
188 double, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t,221 double, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t,
@@ -224,13 +257,10 @@ TEST_F(CannDBDumperUtest, TestCANNDumperShouldReturnTrueWhenComputeTaskDataIsL0T
224 GEInfoDB geInfoDB;257 GEInfoDB geInfoDB;
225 std::string opDescDBPath = Utils::File::PathJoin({TEST_DB_FILE_PATH, geInfoDB.GetDBName()});258 std::string opDescDBPath = Utils::File::PathJoin({TEST_DB_FILE_PATH, geInfoDB.GetDBName()});
226 DBRunner opDescDBRunner(opDescDBPath);259 DBRunner opDescDBRunner(opDescDBPath);
227- std::vector<std::tuple<uint32_t, std::string, uint32_t, uint32_t, uint32_t, uint32_t, std::string, uint32_t,260+ CANNTraceDBDumper::TaskInfoData taskInfoData;
228- std::string, uint32_t, uint32_t, double, uint32_t, uint32_t, std::string,261+ opDescDBRunner.QueryData("select * from TaskInfo", taskInfoData);
229- std::string, std::string, std::string, std::string, std::string, uint32_t,262+ EXPECT_EQ(taskInfoData.size(), 1);
230- uint32_t, uint32_t>> TaskInfoData;263+ EXPECT_EQ(std::get<INPUT_DATA_TYPE_POSITION>(taskInfoData[0]), NA);
231- opDescDBRunner.QueryData("select * from TaskInfo", TaskInfoData);
232- EXPECT_EQ(TaskInfoData.size(), 1);
233- EXPECT_EQ(std::get<INPUT_DATA_TYPE_POSITION>(TaskInfoData[0]), "N/A");
234 264 
235 std::vector<std::tuple<uint32_t, uint32_t, std::string, std::string, uint32_t, std::string,265 std::vector<std::tuple<uint32_t, uint32_t, std::string, std::string, uint32_t, std::string,
236 double, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t,266 double, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t,
@@ -279,20 +309,17 @@ TEST_F(CannDBDumperUtest,
279 GEInfoDB geInfoDB;309 GEInfoDB geInfoDB;
280 std::string opDescDBPath = Utils::File::PathJoin({TEST_DB_FILE_PATH, geInfoDB.GetDBName()});310 std::string opDescDBPath = Utils::File::PathJoin({TEST_DB_FILE_PATH, geInfoDB.GetDBName()});
281 DBRunner opDescDBRunner(opDescDBPath);311 DBRunner opDescDBRunner(opDescDBPath);
282- std::vector<std::tuple<uint32_t, std::string, uint32_t, uint32_t, uint32_t, uint32_t, std::string, uint32_t,312+ CANNTraceDBDumper::TaskInfoData taskInfoData;
283- std::string, uint32_t, uint32_t, double, uint32_t, uint32_t, std::string,313+ opDescDBRunner.QueryData("select * from TaskInfo", taskInfoData);
284- std::string, std::string, std::string, std::string, std::string, uint32_t,314+ EXPECT_EQ(taskInfoData.size(), 1);
285- uint32_t, uint32_t>> TaskInfoData;315+ EXPECT_EQ(std::get<INPUT_DATA_TYPE_POSITION>(taskInfoData[0]), NA);
286- opDescDBRunner.QueryData("select * from TaskInfo", TaskInfoData);
287- EXPECT_EQ(TaskInfoData.size(), 1);
288- EXPECT_EQ(std::get<INPUT_DATA_TYPE_POSITION>(TaskInfoData[0]), "N/A");
289 316 
290 std::vector<std::tuple<uint32_t, uint32_t, std::string, std::string, uint32_t, std::string,317 std::vector<std::tuple<uint32_t, uint32_t, std::string, std::string, uint32_t, std::string,
291 double, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t,318 double, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t, uint32_t,
292 std::string, double, std::string, std::string, uint64_t, std::string, int32_t>> HCCLTaskData;319 std::string, double, std::string, std::string, uint64_t, std::string, int32_t>> HCCLTaskData;
293 hcclOpDBRunner.QueryData("select * from HCCLTask", HCCLTaskData);320 hcclOpDBRunner.QueryData("select * from HCCLTask", HCCLTaskData);
294 EXPECT_EQ(HCCLTaskData.size(), 1);321 EXPECT_EQ(HCCLTaskData.size(), 1);
295- EXPECT_EQ(std::get<GROUP_NAME_POSITION>(HCCLTaskData[0]), "N/A");322+ EXPECT_EQ(std::get<GROUP_NAME_POSITION>(HCCLTaskData[0]), NA);
296}323}
297 324 
298TEST_F(CannDBDumperUtest, TestCANNDumperShouldReturnFalseWhenInsertDataToDBFailed)325TEST_F(CannDBDumperUtest, TestCANNDumperShouldReturnFalseWhenInsertDataToDBFailed)
@@ -385,11 +412,13 @@ TEST_F(CannDBDumperUtest, TestGetFormat_SubFormat_One)
385 EXPECT_EQ(result, "NCHW:1");412 EXPECT_EQ(result, "NCHW:1");
386}413}
387 414 
388-TEST_F(CannDBDumperUtest, TestAddTaskInfoOpIsNullWhenProfLevel0)415+TEST_F(CannDBDumperUtest, TestAddTaskInfoWhenTypeIsReservedAndProfLevel0)
389{416{
390 CANNTraceDBDumper cannTraceDbDumper(TEST_DB_FILE_PATH);417 CANNTraceDBDumper cannTraceDbDumper(TEST_DB_FILE_PATH);
391 auto hostTaskPtr = std::make_shared<HostTask>();418 auto hostTaskPtr = std::make_shared<HostTask>();
392- hostTaskPtr->op = nullptr;419+ auto kernelDesc = std::make_shared<OpDesc>();
420+ auto kernelOp = std::make_shared<Operator>(kernelDesc, 0, OpType::OPTYPE_RESERVED);
421+ hostTaskPtr->op = kernelOp;
393 hostTaskPtr->kernelName = 0;422 hostTaskPtr->kernelName = 0;
394 CANNTraceDBDumper::TaskInfoData taskInfoData;423 CANNTraceDBDumper::TaskInfoData taskInfoData;
395 bool isLevel0 = true;424 bool isLevel0 = true;
@@ -400,11 +429,13 @@ TEST_F(CannDBDumperUtest, TestAddTaskInfoOpIsNullWhenProfLevel0)
400 EXPECT_EQ(taskInfoData.size(), 1ul);429 EXPECT_EQ(taskInfoData.size(), 1ul);
401}430}
402 431 
403-TEST_F(CannDBDumperUtest, TestAddTaskInfoOpIsNullWhenProfLevel1)432+TEST_F(CannDBDumperUtest, TestAddTaskInfoWhenTypeIsReservedAndProfLevel1)
404{433{
405 CANNTraceDBDumper cannTraceDbDumper(TEST_DB_FILE_PATH);434 CANNTraceDBDumper cannTraceDbDumper(TEST_DB_FILE_PATH);
406 auto hostTaskPtr = std::make_shared<HostTask>();435 auto hostTaskPtr = std::make_shared<HostTask>();
407- hostTaskPtr->op = nullptr;436+ auto kernelDesc = std::make_shared<OpDesc>();
437+ auto kernelOp = std::make_shared<Operator>(kernelDesc, 0, OpType::OPTYPE_RESERVED);
438+ hostTaskPtr->op = kernelOp;
408 hostTaskPtr->kernelName = 0;439 hostTaskPtr->kernelName = 0;
409 CANNTraceDBDumper::TaskInfoData taskInfoData;440 CANNTraceDBDumper::TaskInfoData taskInfoData;
410 bool isLevel0 = false;441 bool isLevel0 = false;
@@ -441,4 +472,31 @@ TEST_F(CannDBDumperUtest, TestDumpGeFusionOps)
441 cannTraceDbDumper.DumpGeFusionOps(geFusionOpInfos);472 cannTraceDbDumper.DumpGeFusionOps(geFusionOpInfos);
442 EXPECT_FALSE(cannTraceDbDumper.result_);473 EXPECT_FALSE(cannTraceDbDumper.result_);
443 MOCKER_CPP(&DBRunner::InsertData<>).reset();474 MOCKER_CPP(&DBRunner::InsertData<>).reset();
475+}
476+ 
477+TEST_F(CannDBDumperUtest, TestCANNDumperShouldReturnTrueWhenTaskHasSimtRuntimeTrack)
478+{
479+ MOCKER_CPP(&Analysis::Domain::Environment::Context::IsLevel0).stubs().will(returnValue(false));
480+ MOCKER_CPP(&TypeData::Get).stubs()
481+ .will(returnValue(std::string("KERNEL_SIMT")))
482+ .then(returnValue(std::string("KERNEL_AICORE")));
483+ MockGetComputeTasksWithMixedRuntimeTracks();
484+ CANNTraceDBDumper cannTraceDbDumper(".");
485+ auto treeNode = std::make_shared<TreeNode>(nullptr);
486+ TreeAnalyzer treeAnalyzer(treeNode, THREAD_ID);
487+ bool ret = cannTraceDbDumper.DumpData(treeAnalyzer);
488+ EXPECT_TRUE(ret);
489+ 
490+ GEInfoDB geInfoDB;
491+ std::string opDescDBPath = Utils::File::PathJoin({TEST_DB_FILE_PATH, geInfoDB.GetDBName()});
492+ DBRunner opDescDBRunner(opDescDBPath);
493+ CANNTraceDBDumper::TaskInfoData taskInfoData;
494+ opDescDBRunner.QueryData("select * from TaskInfo", taskInfoData);
495+ EXPECT_EQ(taskInfoData.size(), 2);
496+ EXPECT_EQ(std::get<24>(taskInfoData[0]), "2,3,4");
497+ EXPECT_EQ(std::get<25>(taskInfoData[0]), "5,6,7");
498+ EXPECT_EQ(std::get<4>(taskInfoData[1]), 100u);
499+ EXPECT_EQ(std::get<5>(taskInfoData[1]), 300u);
500+ EXPECT_EQ(std::get<24>(taskInfoData[1]), "N/A");
501+ EXPECT_EQ(std::get<25>(taskInfoData[1]), "N/A");
444}502}
Mtest/msprof_python/ut/testcase/mscalculate/tiling_block_num/test_block_num_calculator.py+1-1
@@ -37,7 +37,7 @@ class TestKfcCalculator(unittest.TestCase):
37 def construct_ge_info_db(self):37 def construct_ge_info_db(self):
38 create_sql = "CREATE TABLE IF NOT EXISTS " + DBNameConstant.TABLE_GE_TASK + \38 create_sql = "CREATE TABLE IF NOT EXISTS " + DBNameConstant.TABLE_GE_TASK + \
39 "(model_id, op_name, stream_id, task_id, block_num, mix_block_num," \39 "(model_id, op_name, stream_id, task_id, block_num, mix_block_num," \
40- "op_stat, task_type, op_type, index_id, thread_id, timestamp, batch_id," \40+ "op_state, task_type, op_type, index_id, thread_id, timestamp, batch_id," \
41 "tensor_num, input_formats, input_data_types, input_shapes," \41 "tensor_num, input_formats, input_data_types, input_shapes," \
42 "output_formats, output_data_types, output_shapes, device_id, context_id, op_flag, hashid)"42 "output_formats, output_data_types, output_shapes, device_id, context_id, op_flag, hashid)"
43 data = (43 data = (
Mtest/st/utils/table_fields.py+253-133
@@ -14,176 +14,296 @@
14# See the Mulan PSL v2 for more details.14# See the Mulan PSL v2 for more details.
15# -------------------------------------------------------------------------15# -------------------------------------------------------------------------
16 16 
17+ 
17class TableFields:18class TableFields:
18- ACC_PMU = [19+ ACC_PMU = ["accId", "readBwLevel", "writeBwLevel", "readOstLevel", "writeOstLevel", "timestampNs", "deviceId"]
19- "accId", "readBwLevel", "writeBwLevel", "readOstLevel", "writeOstLevel", "timestampNs", "deviceId"20+ AICORE_FREQ = ["deviceId", "timestampNs", "freq"]
20- ]21+ CANN_API = ["startNs", "endNs", "type", "globalTid", "connectionId", "name"]
21- AICORE_FREQ = [
22- "deviceId", "timestampNs", "freq"
23- ]
24- CANN_API = [
25- "startNs", "endNs", "type", "globalTid", "connectionId", "name"
26- ]
27 COMMUNICATION_OP = [22 COMMUNICATION_OP = [
28- "opName", "startNs", "endNs", "connectionId", "groupName", "opId", "relay", "retry", "dataType", "algType",23+ "opName",
29- "count", "opType", "deviceId"24+ "startNs",
25+ "endNs",
26+ "connectionId",
27+ "groupName",
28+ "opId",
29+ "relay",
30+ "retry",
31+ "dataType",
32+ "algType",
33+ "count",
34+ "opType",
35+ "deviceId",
30 ]36 ]
31 COMMUNICATION_TASK_INFO = [37 COMMUNICATION_TASK_INFO = [
32- "name", "globalTaskId", "taskType", "planeId", "groupName", "notifyId", "rdmaType", "srcRank", "dstRank",38+ "name",
39+ "globalTaskId",
40+ "taskType",
41+ "planeId",
42+ "groupName",
43+ "notifyId",
44+ "rdmaType",
45+ "srcRank",
46+ "dstRank",
33 "transportType",47 "transportType",
34- "size", "dataType", "linkType", "opId", "isMaster", "bandwidth"48+ "size",
49+ "dataType",
50+ "linkType",
51+ "opId",
52+ "isMaster",
53+ "bandwidth",
35 ]54 ]
36 COMPUTE_TASK_INFO = [55 COMPUTE_TASK_INFO = [
37- "name", "globalTaskId", "blockNum", "mixBlockNum", "taskType", "opType", "inputFormats", "inputDataTypes",56+ "name",
38- "inputShapes", "outputFormats", "outputDataTypes", "outputShapes", "attrInfo", "opState", "hf32Eligible"57+ "globalTaskId",
39- ]58+ "blockNum",
40- CONNECTION_IDS = [59+ "mixBlockNum",
41- "id", "connectionId"60+ "taskType",
42- ]61+ "opType",
43- CPU_USAGE = [62+ "inputFormats",
44- "timestampNs", "cpuId", "usage"63+ "inputDataTypes",
45- ]64+ "inputShapes",
46- ENUM_API_TYPE = [65+ "outputFormats",
47- "id", "name"66+ "outputDataTypes",
48- ]67+ "outputShapes",
49- ENUM_HCCL_DATA_TYPE = [68+ "attrInfo",
50- "id", "name"69+ "opState",
51- ]70+ "hf32Eligible",
52- ENUM_HCCL_LINK_TYPE = [71+ "gridDim",
53- "id", "name"72+ "blockDim",
54- ]
55- ENUM_HCCL_RDMA_TYPE = [
56- "id", "name"
57- ]
58- ENUM_HCCL_TRANSPORT_TYPE = [
59- "id", "name"
60- ]
61- ENUM_MEMCPY_OPERATION = [
62- "id", "name"
63- ]
64- ENUM_MODULE = [
65- "id", "name"
66- ]
67- ENUM_MSTX_EVENT_TYPE = [
68- "id", "name"
69- ]
70- HBM = [
71- "deviceId", "timestampNs", "bandwidth", "hbmId", "type"
72- ]
73- HCCS = [
74- "deviceId", "timestampNs", "txThroughput", "rxThroughput"
75- ]
76- HOST_INFO = [
77- "hostUid", "hostName"
78- ]
79- HOST_MEM_USAGE = [
80- "timestampNs", "usage"
81- ]
82- LLC = [
83- "deviceId", "llcId", "timestampNs", "hitRate", "throughput", "mode"
84- ]
85- MEMCPY_INFO = [
86- "globalTaskId", "size", "memcpyOperation"
87 ]73 ]
74+ CONNECTION_IDS = ["id", "connectionId"]
75+ CPU_USAGE = ["timestampNs", "cpuId", "usage"]
76+ ENUM_API_TYPE = ["id", "name"]
77+ ENUM_HCCL_DATA_TYPE = ["id", "name"]
78+ ENUM_HCCL_LINK_TYPE = ["id", "name"]
79+ ENUM_HCCL_RDMA_TYPE = ["id", "name"]
80+ ENUM_HCCL_TRANSPORT_TYPE = ["id", "name"]
81+ ENUM_MEMCPY_OPERATION = ["id", "name"]
82+ ENUM_MODULE = ["id", "name"]
83+ ENUM_MSTX_EVENT_TYPE = ["id", "name"]
84+ HBM = ["deviceId", "timestampNs", "bandwidth", "hbmId", "type"]
85+ HCCS = ["deviceId", "timestampNs", "txThroughput", "rxThroughput"]
86+ HOST_INFO = ["hostUid", "hostName"]
87+ HOST_MEM_USAGE = ["timestampNs", "usage"]
88+ LLC = ["deviceId", "llcId", "timestampNs", "hitRate", "throughput", "mode"]
89+ MEMCPY_INFO = ["globalTaskId", "size", "memcpyOperation"]
88 MEMORY_RECORD = [90 MEMORY_RECORD = [
89- "component", "timestamp", "totalAllocated", "totalReserved", "totalActive", "streamPtr", "deviceId"91+ "component",
90- ]92+ "timestamp",
91- META_DATA = [93+ "totalAllocated",
92- "name", "value"94+ "totalReserved",
95+ "totalActive",
96+ "streamPtr",
97+ "deviceId",
93 ]98 ]
99+ META_DATA = ["name", "value"]
94 NETDEV_STATS = [100 NETDEV_STATS = [
95- "deviceId", "timestampNs", "macTxPfcPkt", "macRxPfcPkt", "macTxByte", "macTxBandwidth", "macRxByte",101+ "deviceId",
102+ "timestampNs",
103+ "macTxPfcPkt",
104+ "macRxPfcPkt",
105+ "macTxByte",
106+ "macTxBandwidth",
107+ "macRxByte",
96 "macRxBandwidth",108 "macRxBandwidth",
97- "macTxBadByte", "macRxBadByte", "roceTxPkt", "roceRxPkt", "roceTxErrPkt", "roceRxErrPkt", "roceTxCnpPkt",109+ "macTxBadByte",
110+ "macRxBadByte",
111+ "roceTxPkt",
112+ "roceRxPkt",
113+ "roceTxErrPkt",
114+ "roceRxErrPkt",
115+ "roceTxCnpPkt",
98 "roceRxCnpPkt",116 "roceRxCnpPkt",
99- "roceNewPktRty", "nicTxByte", "nicTxBandwidth", "nicRxByte", "nicRxBandwidth"117+ "roceNewPktRty",
118+ "nicTxByte",
119+ "nicTxBandwidth",
120+ "nicRxByte",
121+ "nicRxBandwidth",
100 ]122 ]
101 NIC = [123 NIC = [
102- "deviceId", "timestampNs", "bandwidth", "rxPacketRate", "rxByteRate", "rxPackets", "rxBytes", "rxErrors",124+ "deviceId",
125+ "timestampNs",
126+ "bandwidth",
127+ "rxPacketRate",
128+ "rxByteRate",
129+ "rxPackets",
130+ "rxBytes",
131+ "rxErrors",
103 "rxDropped",132 "rxDropped",
104- "txPacketRate", "txByteRate", "txPackets", "txBytes", "txErrors", "txDropped", "funcId"133+ "txPacketRate",
105- ]134+ "txByteRate",
106- NPU_INFO = [135+ "txPackets",
107- "id", "name"136+ "txBytes",
108- ]137+ "txErrors",
109- NPU_MEM = [138+ "txDropped",
110- "type", "ddr", "hbm", "timestampNs", "deviceId"139+ "funcId",
111- ]
112- NPU_MODULE_MEM = [
113- "moduleId", "timestampNs", "totalReserved", "deviceId"
114 ]140 ]
141+ NPU_INFO = ["id", "name"]
142+ NPU_MEM = ["type", "ddr", "hbm", "timestampNs", "deviceId"]
143+ NPU_MODULE_MEM = ["moduleId", "timestampNs", "totalReserved", "deviceId"]
115 NPU_OP_MEM = [144 NPU_OP_MEM = [
116- "operatorName", "addr", "type", "size", "timestampNs", "globalTid", "totalAllocate", "totalReserve",145+ "operatorName",
146+ "addr",
147+ "type",
148+ "size",
149+ "timestampNs",
150+ "globalTid",
151+ "totalAllocate",
152+ "totalReserve",
117 "component",153 "component",
118- "deviceId"154+ "deviceId",
119 ]155 ]
120 OP_MEMORY = [156 OP_MEMORY = [
121- "name", "size", "allocationTime", "releaseTime", "activeReleaseTime", "duration", "activeDuration",157+ "name",
158+ "size",
159+ "allocationTime",
160+ "releaseTime",
161+ "activeReleaseTime",
162+ "duration",
163+ "activeDuration",
122 "allocationTotalAllocated",164 "allocationTotalAllocated",
123- "allocationTotalReserved", "allocationTotalActive", "releaseTotalAllocated", "releaseTotalReserved",165+ "allocationTotalReserved",
166+ "allocationTotalActive",
167+ "releaseTotalAllocated",
168+ "releaseTotalReserved",
124 "releaseTotalActive",169 "releaseTotalActive",
125- "streamPtr", "deviceId"170+ "streamPtr",
171+ "deviceId",
126 ]172 ]
127 PCIE = [173 PCIE = [
128- "deviceId", "timestampNs", "txPostMin", "txPostMax", "txPostAvg", "txNonpostMin", "txNonpostMax",174+ "deviceId",
129- "txNonpostAvg", "txCplMin",175+ "timestampNs",
130- "txCplMax", "txCplAvg", "txNonpostLatencyMin", "txNonpostLatencyMax", "txNonpostLatencyAvg", "rxPostMin",176+ "txPostMin",
177+ "txPostMax",
178+ "txPostAvg",
179+ "txNonpostMin",
180+ "txNonpostMax",
181+ "txNonpostAvg",
182+ "txCplMin",
183+ "txCplMax",
184+ "txCplAvg",
185+ "txNonpostLatencyMin",
186+ "txNonpostLatencyMax",
187+ "txNonpostLatencyAvg",
188+ "rxPostMin",
131 "rxPostMax",189 "rxPostMax",
132- "rxPostAvg", "rxNonpostMin", "rxNonpostMax", "rxNonpostAvg", "rxCplMin", "rxCplMax", "rxCplAvg"190+ "rxPostAvg",
191+ "rxNonpostMin",
192+ "rxNonpostMax",
193+ "rxNonpostAvg",
194+ "rxCplMin",
195+ "rxCplMax",
196+ "rxCplAvg",
133 ]197 ]
134 PYTORCH_API = [198 PYTORCH_API = [
135- "startNs", "endNs", "globalTid", "connectionId", "name", "sequenceNumber", "fwdThreadId", "inputDtypes",199+ "startNs",
200+ "endNs",
201+ "globalTid",
202+ "connectionId",
203+ "name",
204+ "sequenceNumber",
205+ "fwdThreadId",
206+ "inputDtypes",
136 "inputShapes",207 "inputShapes",
137- "callchainId", "type"208+ "callchainId",
138- ]209+ "type",
139- PYTORCH_CALLCHAINS = [
140- "id", "stack", "stackDepth"
141- ]
142- QOS = [
143- "deviceId", "eventName", "bandwidth", "timestampNs"
144 ]210 ]
211+ PYTORCH_CALLCHAINS = ["id", "stack", "stackDepth"]
212+ QOS = ["deviceId", "eventName", "bandwidth", "timestampNs"]
145 ROCE = [213 ROCE = [
146- "deviceId", "timestampNs", "bandwidth", "rxPacketRate", "rxByteRate", "rxPackets", "rxBytes", "rxErrors",214+ "deviceId",
215+ "timestampNs",
216+ "bandwidth",
217+ "rxPacketRate",
218+ "rxByteRate",
219+ "rxPackets",
220+ "rxBytes",
221+ "rxErrors",
147 "rxDropped",222 "rxDropped",
148- "txPacketRate", "txByteRate", "txPackets", "txBytes", "txErrors", "txDropped", "funcId"223+ "txPacketRate",
149- ]224+ "txByteRate",
150- SESSION_TIME_INFO = [225+ "txPackets",
151- "startTimeNs", "endTimeNs"226+ "txBytes",
152- ]227+ "txErrors",
153- SOC_BANDWIDTH_LEVEL = [228+ "txDropped",
154- "l2BufferBwLevel", "mataBwLevel", "timestampNs", "deviceId"229+ "funcId",
155- ]
156- STEP_TIME = [
157- "id", "startNs", "endNs"
158- ]
159- STRING_IDS = [
160- "id", "value"
161 ]230 ]
231+ SESSION_TIME_INFO = ["startTimeNs", "endTimeNs"]
232+ SOC_BANDWIDTH_LEVEL = ["l2BufferBwLevel", "mataBwLevel", "timestampNs", "deviceId"]
233+ STEP_TIME = ["id", "startNs", "endNs"]
234+ STRING_IDS = ["id", "value"]
162 TASK = [235 TASK = [
163- "startNs", "endNs", "deviceId", "connectionId", "globalTaskId", "globalPid", "taskType", "contextId",236+ "startNs",
164- "streamId", "taskId", "modelId"237+ "endNs",
238+ "deviceId",
239+ "connectionId",
240+ "globalTaskId",
241+ "globalPid",
242+ "taskType",
243+ "contextId",
244+ "streamId",
245+ "taskId",
246+ "modelId",
165 ]247 ]
166- TASK_PMU_INFO = [248+ TASK_PMU_INFO = ["globalTaskId", "name", "value"]
167- "globalTaskId", "name", "value"249+ RANK_DEVICE_MAP = ["rankId", "deviceId"]
250+ ClusterCommunicationBandwidth = [
251+ "step",
252+ "rank_id",
253+ "hccl_op_name",
254+ "group_name",
255+ "band_type",
256+ "transit_size",
257+ "transit_time",
258+ "bandwidth",
259+ "large_packet_ratio",
260+ "package_size",
261+ "count",
262+ "total_duration",
168 ]263 ]
169- RANK_DEVICE_MAP = [264+ ClusterCommunicationMatrix = [
170- "rankId", "deviceId"265+ "step",
266+ "hccl_op_name",
267+ "group_name",
268+ "src_rank",
269+ "dst_rank",
270+ "transport_type",
271+ "op_name",
272+ "transit_size",
273+ "transit_time",
274+ "bandwidth",
275+ ]
276+ ClusterCommunicationTime = [
277+ "step",
278+ "rank_id",
279+ "hccl_op_name",
280+ "group_name",
281+ "start_timestamp",
282+ "elapsed_time",
283+ "transit_time",
284+ "wait_time",
285+ "synchronization_time",
286+ "idle_time",
287+ "synchronization_time_ratio",
288+ "wait_time_ratio",
289+ ]
290+ ClusterStepTraceTime = [
291+ "step",
292+ "type",
293+ "index",
294+ "computing",
295+ "communication_not_overlapped",
296+ "overlapped",
297+ "communication",
298+ "free",
299+ "stage",
300+ "bubble",
301+ "communication_not_overlapped_and_exclude_receive",
302+ "preparing",
303+ "dp_index",
304+ "pp_index",
305+ "tp_index",
171 ]306 ]
172- ClusterCommunicationBandwidth = ["step", "rank_id", "hccl_op_name", "group_name", "band_type", "transit_size",
173- "transit_time", "bandwidth", "large_packet_ratio", "package_size", "count",
174- "total_duration"]
175- ClusterCommunicationMatrix = ["step", "hccl_op_name", "group_name", "src_rank", "dst_rank", "transport_type",
176- "op_name",
177- "transit_size", "transit_time", "bandwidth"]
178- ClusterCommunicationTime = ["step", "rank_id", "hccl_op_name", "group_name", "start_timestamp", "elapsed_time",
179- "transit_time",
180- "wait_time", "synchronization_time", "idle_time", "synchronization_time_ratio",
181- "wait_time_ratio"]
182- ClusterStepTraceTime = ["step", "type", "index", "computing", "communication_not_overlapped", "overlapped",
183- "communication",
184- "free", "stage", "bubble", "communication_not_overlapped_and_exclude_receive", "preparing",
185- "dp_index",
186- "pp_index", "tp_index"]
187 CommunicationGroupMapping = ["type", "rank_set", "group_name", "group_id", "pg_name"]307 CommunicationGroupMapping = ["type", "rank_set", "group_name", "group_id", "pg_name"]
188 HostInfo = ["hostUid", "hostName"]308 HostInfo = ["hostUid", "hostName"]
189 RankDeviceMap = ["rankId", "deviceId", "hostUid", "profilePath"]309 RankDeviceMap = ["rankId", "deviceId", "hostUid", "profilePath"]