AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

SWDEV-375202 - Fixes print formatting for flags in hipGraphDebugPrint() to match CUDA

Browse Source 
Change-Id: I2d85fc38d2c65bc12534109883fe00802e77e62d
This commit is contained in:
Sourabh Betigeri
2023-02-08 19:28:51 +00:00
committed by Sourabh Betigeri
parent 810c4ac244
commit 4cbf9a4a16
1 changed files with 44 additions and 4 deletions
Download Patch File Download Diff File Expand all files Collapse all files
hipamd/src/hip_graph_internal.hpp
+44 -4
View File
@@ -735,12 +735,31 @@ class hipGraphKernelNode : public hipGraphNode {
unsigned int kernelAttrInUse_;
public:
void PrintAttributes(std::ostream& out, hipGraphDebugDotFlags flag) {
out << "[";
out << "style";
out << "=\"";
out << style_;
(flag == hipGraphDebugDotFlagsKernelNodeParams ||
flag == hipGraphDebugDotFlagsKernelNodeAttributes) ?
out << "\n" : out << "\"";
out << "shape";
out << "=\"";
out << GetShape(flag);
out << "\"";
out << "label";
out << "=\"";
out << GetLabel(flag);
out << "\"";
out << "];";
}
std::string GetLabel(hipGraphDebugDotFlags flag) {
hipFunction_t func = getFunc(*pKernelParams_, ihipGetDevice());
hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func);
std::string label;
if (flag == hipGraphDebugDotFlagsKernelNodeParams || flag == hipGraphDebugDotFlagsVerbose) {
char buffer[500];
char buffer[500];
if (flag == hipGraphDebugDotFlagsVerbose) {
sprintf(buffer,
"{\n%s\n| {ID | %d | %s\\<\\<\\<(%u,%u,%u),(%u,%u,%u),%u\\>\\>\\>}\n| {{node "
"handle | func handle} | {%p | %p}}\n| {accessPolicyWindow | {base_ptr | num_bytes | "
@@ -754,8 +773,29 @@ class hipGraphKernelNode : public hipGraphNode {
kernelAttr_.accessPolicyWindow.hitRatio, kernelAttr_.accessPolicyWindow.hitProp,
kernelAttr_.accessPolicyWindow.missProp, kernelAttr_.cooperative);
label = buffer;
} else {
label = std::to_string(GetID()) + "\n" + function->name();
}
else if (flag == hipGraphDebugDotFlagsKernelNodeAttributes) {
sprintf(buffer,
"{\n%s\n| {ID | %d | %s}\n"
"| {accessPolicyWindow | {base_ptr | num_bytes | "
"hitRatio | hitProp | missProp} |\n| {%p | %ld | %f | %d | %d}}\n| {cooperative | "
"%u}\n| {priority | 0}\n}",
label_.c_str(), GetID(), function->name().c_str(),
kernelAttr_.accessPolicyWindow.base_ptr, kernelAttr_.accessPolicyWindow.num_bytes,
kernelAttr_.accessPolicyWindow.hitRatio, kernelAttr_.accessPolicyWindow.hitProp,
kernelAttr_.accessPolicyWindow.missProp, kernelAttr_.cooperative);
label = buffer;
}
else if (flag == hipGraphDebugDotFlagsKernelNodeParams) {
sprintf(buffer, "%d\n%s\n\\<\\<\\<(%u,%u,%u),(%u,%u,%u),%u\\>\\>\\>",
GetID(), function->name().c_str(), pKernelParams_->gridDim.x,
pKernelParams_->gridDim.y, pKernelParams_->gridDim.z,
pKernelParams_->blockDim.x, pKernelParams_->blockDim.y,
pKernelParams_->blockDim.z, pKernelParams_->sharedMemBytes);
label = buffer;
}
else {
label = std::to_string(GetID()) + "\n" + function->name() + "\n";
}
return label;
}