AI/rocm-systems
2
0
Tükrözés 0
Kód Hibajegyek Egyesítési kérések Actions Packages Projektek Kiadások Wiki Tevékenység

SWDEV-325711 - Correct formatting

Forráskód böngészése 
Change-Id: Ie26159e0bb3315cf7c3de1eb682f23ef343df0f2


[ROCm/clr commit: 16b31b0c54]
This commit is contained in:
Anusha GodavarthySurya
2022-11-30 05:15:01 +00:00
szülő 5780c7066f
commit e7fa085afe
3 fájl változott, egészen pontosan 92 új sor hozzáadva és 89 régi sor törölve
Download Patch File Download Diff File Expand all files Collapse all files
projects/clr/hipamd/src/hip_graph.cpp
+38 -30
Fájl megtekintése
@@ -853,8 +853,7 @@ hipError_t hipStreamBeginCapture_common(hipStream_t stream, hipStreamCaptureMode
if (stream == nullptr) {
return hipErrorStreamCaptureUnsupported;
}
if (mode < hipStreamCaptureModeGlobal ||
mode > hipStreamCaptureModeRelaxed) {
if (mode < hipStreamCaptureModeGlobal || mode > hipStreamCaptureModeRelaxed) {
return hipErrorInvalidValue;
}
hip::Stream* s = reinterpret_cast<hip::Stream*>(stream);
@@ -909,8 +908,7 @@ hipError_t hipStreamEndCapture_common(hipStream_t stream, hipGraph_t* pGraph) {
}
// If mode is not hipStreamCaptureModeRelaxed, hipStreamEndCapture must be called on the stream
// from the same thread
const auto& it = std::find(hip::tls.capture_streams_.begin(),
hip::tls.capture_streams_.end(), s);
const auto& it = std::find(hip::tls.capture_streams_.begin(), hip::tls.capture_streams_.end(), s);
if (s->GetCaptureMode() != hipStreamCaptureModeRelaxed) {
if (it == hip::tls.capture_streams_.end()) {
return hipErrorStreamCaptureWrongThread;
@@ -1118,7 +1116,8 @@ hipError_t ihipGraphInstantiate(hipGraphExec_t* pGraphExec, hipGraph_t graph) {
std::vector<Node> levelOrder;
clonedGraph->LevelOrder(levelOrder);
clonedGraph->GetUserObjs(graphExeUserObj);
*pGraphExec = new hipGraphExec(levelOrder, parallelLists, nodeWaitLists, clonedNodes, graphExeUserObj);
*pGraphExec =
new hipGraphExec(levelOrder, parallelLists, nodeWaitLists, clonedNodes, graphExeUserObj);
if (*pGraphExec != nullptr) {
return (*pGraphExec)->Init();
} else {
@@ -1167,7 +1166,7 @@ hipError_t ihipGraphLaunch(hipGraphExec_t graphExec, hipStream_t stream) {
return graphExec->Run(stream);
}
hipError_t hipGraphLaunch_common( hipGraphExec_t graphExec, hipStream_t stream) {
hipError_t hipGraphLaunch_common(hipGraphExec_t graphExec, hipStream_t stream) {
if (graphExec == nullptr || !hip::isValid(stream) || !hipGraphExec::isGraphExecValid(graphExec)) {
return hipErrorInvalidValue;
}
@@ -1270,7 +1269,8 @@ hipError_t hipGraphKernelNodeSetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt
if (hNode == nullptr || value == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
if (attr != hipKernelNodeAttributeAccessPolicyWindow && attr != hipKernelNodeAttributeCooperative ) {
if (attr != hipKernelNodeAttributeAccessPolicyWindow &&
attr != hipKernelNodeAttributeCooperative) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphKernelNode*>(hNode)->SetAttrParams(attr, value));
@@ -1282,7 +1282,8 @@ hipError_t hipGraphKernelNodeGetAttribute(hipGraphNode_t hNode, hipKernelNodeAtt
if (hNode == nullptr || value == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
if (attr != hipKernelNodeAttributeAccessPolicyWindow && attr != hipKernelNodeAttributeCooperative ) {
if (attr != hipKernelNodeAttributeAccessPolicyWindow &&
attr != hipKernelNodeAttributeCooperative) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphKernelNode*>(hNode)->GetAttrParams(attr, value));
@@ -1423,7 +1424,7 @@ hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGra
HIP_RETURN(hipErrorUnknown);
}
// Validate whether the topology of node and childGraph matches
// Validate whether the topology of node and childGraph matches
std::vector<Node> childGraphNodes1;
node->LevelOrder(childGraphNodes1);
@@ -1449,8 +1450,9 @@ hipError_t hipGraphExecChildGraphNodeSetParams(hipGraphExec_t hGraphExec, hipGra
HIP_RETURN(reinterpret_cast<hipChildGraphNode*>(clonedNode)->SetParams(childGraph));
}
hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
unsigned long long* pId) {
hipError_t hipStreamGetCaptureInfo_common(hipStream_t stream,
hipStreamCaptureStatus* pCaptureStatus,
unsigned long long* pId) {
if (pCaptureStatus == nullptr || !hip::isValid(stream)) {
return hipErrorInvalidValue;
}
@@ -1472,16 +1474,17 @@ hipError_t hipStreamGetCaptureInfo(hipStream_t stream, hipStreamCaptureStatus* p
}
hipError_t hipStreamGetCaptureInfo_spt(hipStream_t stream, hipStreamCaptureStatus* pCaptureStatus,
unsigned long long* pId) {
unsigned long long* pId) {
HIP_INIT_API(hipStreamGetCaptureInfo, stream, pCaptureStatus, pId);
PER_THREAD_DEFAULT_STREAM(stream);
HIP_RETURN(hipStreamGetCaptureInfo_common(stream, pCaptureStatus, pId));
}
hipError_t hipStreamGetCaptureInfo_v2_common(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
unsigned long long* id_out, hipGraph_t* graph_out,
const hipGraphNode_t** dependencies_out,
size_t* numDependencies_out) {
hipError_t hipStreamGetCaptureInfo_v2_common(hipStream_t stream,
hipStreamCaptureStatus* captureStatus_out,
unsigned long long* id_out, hipGraph_t* graph_out,
const hipGraphNode_t** dependencies_out,
size_t* numDependencies_out) {
if (captureStatus_out == nullptr) {
return hipErrorInvalidValue;
}
@@ -1517,18 +1520,19 @@ hipError_t hipStreamGetCaptureInfo_v2(hipStream_t stream, hipStreamCaptureStatus
HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out,
dependencies_out, numDependencies_out);
HIP_RETURN(hipStreamGetCaptureInfo_v2_common(stream, captureStatus_out, id_out, graph_out,
dependencies_out, numDependencies_out));
dependencies_out, numDependencies_out));
}
hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream, hipStreamCaptureStatus* captureStatus_out,
unsigned long long* id_out, hipGraph_t* graph_out,
const hipGraphNode_t** dependencies_out,
size_t* numDependencies_out) {
hipError_t hipStreamGetCaptureInfo_v2_spt(hipStream_t stream,
hipStreamCaptureStatus* captureStatus_out,
unsigned long long* id_out, hipGraph_t* graph_out,
const hipGraphNode_t** dependencies_out,
size_t* numDependencies_out) {
HIP_INIT_API(hipStreamGetCaptureInfo_v2, stream, captureStatus_out, id_out, graph_out,
dependencies_out, numDependencies_out);
PER_THREAD_DEFAULT_STREAM(stream);
HIP_RETURN(hipStreamGetCaptureInfo_v2_common(stream, captureStatus_out, id_out, graph_out,
dependencies_out, numDependencies_out));
dependencies_out, numDependencies_out));
}
hipError_t hipStreamUpdateCaptureDependencies(hipStream_t stream, hipGraphNode_t* dependencies,
@@ -1947,8 +1951,8 @@ hipError_t hipGraphHostNodeGetParams(hipGraphNode_t node, hipHostNodeParams* pNo
hipError_t hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParams* pNodeParams) {
HIP_INIT_API(hipGraphHostNodeSetParams, node, pNodeParams);
if (pNodeParams == nullptr || pNodeParams->fn == nullptr ||
node == nullptr || !hipGraphNode::isNodeValid(node)) {
if (pNodeParams == nullptr || pNodeParams->fn == nullptr || node == nullptr ||
!hipGraphNode::isNodeValid(node)) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphHostNode*>(node)->SetParams(pNodeParams));
@@ -1957,7 +1961,8 @@ hipError_t hipGraphHostNodeSetParams(hipGraphNode_t node, const hipHostNodeParam
hipError_t hipGraphExecHostNodeSetParams(hipGraphExec_t hGraphExec, hipGraphNode_t node,
const hipHostNodeParams* pNodeParams) {
HIP_INIT_API(hipGraphExecHostNodeSetParams, hGraphExec, node, pNodeParams);
if (hGraphExec == nullptr || pNodeParams == nullptr || pNodeParams->fn == nullptr || node == nullptr || !hipGraphNode::isNodeValid(node)) {
if (hGraphExec == nullptr || pNodeParams == nullptr || pNodeParams->fn == nullptr ||
node == nullptr || !hipGraphNode::isNodeValid(node)) {
HIP_RETURN(hipErrorInvalidValue);
}
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(node);
@@ -2061,9 +2066,11 @@ hipError_t hipDeviceGraphMemTrim(int device) {
return HIP_RETURN(hipSuccess);
}
hipError_t hipUserObjectCreate(hipUserObject_t* object_out, void* ptr, hipHostFn_t destroy, unsigned int initialRefcount, unsigned int flags) {
hipError_t hipUserObjectCreate(hipUserObject_t* object_out, void* ptr, hipHostFn_t destroy,
unsigned int initialRefcount, unsigned int flags) {
HIP_INIT_API(hipUserObjectCreate, object_out, ptr, destroy, initialRefcount, flags);
if (object_out == nullptr || flags != hipUserObjectNoDestructorSync || initialRefcount == 0 || destroy == nullptr || initialRefcount > INT_MAX) {
if (object_out == nullptr || flags != hipUserObjectNoDestructorSync || initialRefcount == 0 ||
destroy == nullptr || initialRefcount > INT_MAX) {
HIP_RETURN(hipErrorInvalidValue);
}
@@ -2099,7 +2106,8 @@ hipError_t hipUserObjectRetain(hipUserObject_t object, unsigned int count) {
HIP_RETURN(hipSuccess);
}
hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count, unsigned int flags) {
hipError_t hipGraphRetainUserObject(hipGraph_t graph, hipUserObject_t object, unsigned int count,
unsigned int flags) {
HIP_INIT_API(hipGraphRetainUserObject, graph, object, count, flags);
hipError_t status = hipSuccess;
if (graph == nullptr || object == nullptr || count == 0 || count > INT_MAX ||
@@ -2141,7 +2149,7 @@ hipError_t hipGraphReleaseUserObject(hipGraph_t graph, hipUserObject_t object, u
hipError_t hipGraphKernelNodeCopyAttributes(hipGraphNode_t hSrc, hipGraphNode_t hDst) {
HIP_INIT_API(hipGraphKernelNodeCopyAttributes, hSrc, hDst);
if(hSrc == nullptr || hDst == nullptr) {
if (hSrc == nullptr || hDst == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
HIP_RETURN(reinterpret_cast<hipGraphKernelNode*>(hDst)->CopyAttr(
@@ -2173,7 +2181,7 @@ hipError_t hipGraphDebugDotPrint(hipGraph_t graph, const char* path, unsigned in
hipError_t hipGraphNodeSetEnabled(hipGraphExec_t hGraphExec, hipGraphNode_t hNode,
unsigned int isEnabled) {
HIP_INIT_API(hipGraphNodeSetEnabled, hGraphExec, hNode, isEnabled);
if (hGraphExec == nullptr || hNode == nullptr) {
if (hGraphExec == nullptr || hNode == nullptr) {
HIP_RETURN(hipErrorInvalidValue);
}
hipGraphNode_t clonedNode = hGraphExec->GetClonedNode(hNode);
projects/clr/hipamd/src/hip_graph_internal.cpp
+10 -13
Fájl megtekintése
@@ -37,7 +37,7 @@ const char* GetGraphNodeTypeString(uint32_t op) {
CASE_STRING(hipGraphNodeTypeWaitEvent, WaitEventNode)
CASE_STRING(hipGraphNodeTypeEventRecord, EventRecordNode)
CASE_STRING(hipGraphNodeTypeExtSemaphoreSignal, ExtSemaphoreSignalNode)
CASE_STRING(hipGraphNodeTypeExtSemaphoreWait, ExtSemaphoreWaitNode)
CASE_STRING(hipGraphNodeTypeExtSemaphoreWait, ExtSemaphoreWaitNode)
CASE_STRING(hipGraphNodeTypeMemcpyFromSymbol, MemcpyFromSymbolNode)
CASE_STRING(hipGraphNodeTypeMemcpyToSymbol, MemcpyToSymbolNode)
default:
@@ -73,14 +73,14 @@ hipError_t hipGraphMemcpyNode1D::ValidateParams(void* dst, const void* src, size
size_t dOffset = 0;
amd::Memory* dstMemory = getMemoryObject(dst, dOffset);
if ((srcMemory == nullptr) && (dstMemory != nullptr)) { //host to device
if ((srcMemory == nullptr) && (dstMemory != nullptr)) { // host to device
if (origDstMemory->getContext().devices()[0] != dstMemory->getContext().devices()[0]) {
return hipErrorInvalidValue;
}
if (kind != hipMemcpyHostToDevice) {
return hipErrorInvalidValue;
}
} else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { //device to host
} else if ((srcMemory != nullptr) && (dstMemory == nullptr)) { // device to host
if (origSrcMemory->getContext().devices()[0] != srcMemory->getContext().devices()[0]) {
return hipErrorInvalidValue;
}
@@ -636,14 +636,12 @@ void ihipGraph::LevelOrder(std::vector<Node>& levelOrder) {
}
}
const ihipGraph* ihipGraph::getOriginalGraph() const {
return pOriginalGraph_;
}
const ihipGraph* ihipGraph::getOriginalGraph() const { return pOriginalGraph_; }
void ihipGraph::setOriginalGraph(const ihipGraph* pOriginalGraph) {
pOriginalGraph_ = pOriginalGraph;
}
ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const{
ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const {
ihipGraph* newGraph = new ihipGraph();
for (auto entry : vertices_) {
hipGraphNode* node = entry->clone();
@@ -674,7 +672,7 @@ ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const{
return newGraph;
}
ihipGraph* ihipGraph::clone() const{
ihipGraph* ihipGraph::clone() const {
std::unordered_map<Node, Node> clonedNodes;
return clone(clonedNodes);
}
@@ -691,9 +689,9 @@ hipError_t hipGraphExec::CreateQueues(size_t numQueues) {
parallelQueues_.reserve(numQueues);
for (size_t i = 0; i < numQueues; i++) {
amd::HostQueue* queue;
queue = new amd::HostQueue(*hip::getCurrentDevice()->asContext(),
*hip::getCurrentDevice()->devices()[0], 0,
amd::CommandQueue::RealTimeDisabled, amd::CommandQueue::Priority::Normal);
queue = new amd::HostQueue(
*hip::getCurrentDevice()->asContext(), *hip::getCurrentDevice()->devices()[0], 0,
amd::CommandQueue::RealTimeDisabled, amd::CommandQueue::Priority::Normal);
bool result = (queue != nullptr) ? queue->create() : false;
// Create a host queue
@@ -720,8 +718,7 @@ hipError_t hipGraphExec::Init() {
hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
std::unordered_map<Node, std::vector<Node>>& nodeWaitLists,
std::vector<Node>& levelOrder,
std::vector<amd::Command*>& rootCommands,
std::vector<Node>& levelOrder, std::vector<amd::Command*>& rootCommands,
amd::Command*& endCommand, amd::HostQueue* queue) {
hipError_t status;
for (auto& node : levelOrder) {
projects/clr/hipamd/src/hip_graph_internal.hpp
+44 -46
Fájl megtekintése
@@ -36,19 +36,19 @@
typedef hipGraphNode* Node;
hipError_t FillCommands(std::vector<std::vector<Node>>& parallelLists,
std::unordered_map<Node, std::vector<Node>>& nodeWaitLists,
std::vector<Node>& levelOrder,
std::vector<amd::Command*>& rootCommands,
std::vector<Node>& levelOrder, std::vector<amd::Command*>& rootCommands,
amd::Command*& endCommand, amd::HostQueue* queue);
void UpdateQueue(std::vector<std::vector<Node>>& parallelLists, amd::HostQueue*& queue,
hipGraphExec* ptr);
struct hipUserObject : public amd::ReferenceCountedObject {
typedef void(*UserCallbackDestructor)(void* data);
typedef void (*UserCallbackDestructor)(void* data);
static std::unordered_set<hipUserObject*> ObjectSet_;
static amd::Monitor UserObjectLock_;
public:
hipUserObject(UserCallbackDestructor callback, void* data, unsigned int flags) : ReferenceCountedObject(),
callback_(callback), data_(data), flags_(flags) {
hipUserObject(UserCallbackDestructor callback, void* data, unsigned int flags)
: ReferenceCountedObject(), callback_(callback), data_(data), flags_(flags) {
amd::ScopedLock lock(UserObjectLock_);
ObjectSet_.insert(this);
}
@@ -88,6 +88,7 @@ struct hipUserObject : public amd::ReferenceCountedObject {
ObjectSet_.erase(pUsertObj);
}
}
private:
UserCallbackDestructor callback_;
void* data_;
@@ -431,9 +432,7 @@ struct ihipGraph {
return true;
}
// Delete user obj resource from graph
void RemoveUserObjGraph(hipUserObject* pUserObj) {
graphUserObj_.erase(pUserObj);
}
void RemoveUserObjGraph(hipUserObject* pUserObj) { graphUserObj_.erase(pUserObj); }
// saves the original graph ptr if cloned
void setOriginalGraph(const ihipGraph* pOriginalGraph);
@@ -443,9 +442,8 @@ struct ihipGraph {
void GetRunList(std::vector<std::vector<Node>>& parallelLists,
std::unordered_map<Node, std::vector<Node>>& dependencies);
void LevelOrder(std::vector<Node>& levelOrder);
void GetUserObjs( std::unordered_set<hipUserObject*>& graphExeUserObjs) {
for(auto userObj : graphUserObj_)
{
void GetUserObjs(std::unordered_set<hipUserObject*>& graphExeUserObjs) {
for (auto userObj : graphUserObj_) {
userObj->retain();
graphExeUserObjs.insert(userObj);
}
@@ -458,7 +456,7 @@ struct ihipGraph {
for (auto node : vertices_) {
node->GenerateDOTNode(GetID(), fout, flag);
}
fout<<"\n";
fout << "\n";
for (auto& node : vertices_) {
node->GenerateDOTNodeEdges(GetID(), fout, flag);
}
@@ -485,7 +483,8 @@ struct hipGraphExec {
public:
hipGraphExec(std::vector<Node>& levelOrder, std::vector<std::vector<Node>>& lists,
std::unordered_map<Node, std::vector<Node>>& nodeWaitLists,
std::unordered_map<Node, Node>& clonedNodes, std::unordered_set<hipUserObject*>& userObjs)
std::unordered_map<Node, Node>& clonedNodes,
std::unordered_set<hipUserObject*>& userObjs)
: parallelLists_(lists),
levelOrder_(levelOrder),
nodeWaitLists_(nodeWaitLists),
@@ -713,25 +712,25 @@ class hipGraphKernelNode : public hipGraphNode {
if (!func) {
return hipErrorInvalidDeviceFunction;
}
hip::DeviceFunc *function = hip::DeviceFunc::asFunction(func);
hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func);
amd::Kernel* kernel = function->kernel();
const amd::KernelSignature& signature = kernel->signature();
numParams_ = signature.numParameters();
// Allocate/assign memory if params are passed part of 'kernelParams'
if (pNodeParams->kernelParams != nullptr) {
pKernelParams_->kernelParams = (void**) malloc(numParams_ * sizeof(void*));
pKernelParams_->kernelParams = (void**)malloc(numParams_ * sizeof(void*));
if (pKernelParams_->kernelParams == nullptr) {
return hipErrorOutOfMemory;
}
for (uint32_t i = 0; i < numParams_; ++i) {
const amd::KernelParameterDescriptor& desc = signature.at(i);
pKernelParams_->kernelParams[i] = malloc(desc.size_);
if (pKernelParams_->kernelParams[i] == nullptr) {
return hipErrorOutOfMemory;
}
::memcpy(pKernelParams_->kernelParams[i], (pNodeParams->kernelParams[i]), desc.size_);
pKernelParams_->kernelParams[i] = malloc(desc.size_);
if (pKernelParams_->kernelParams[i] == nullptr) {
return hipErrorOutOfMemory;
}
::memcpy(pKernelParams_->kernelParams[i], (pNodeParams->kernelParams[i]), desc.size_);
}
}
@@ -742,22 +741,22 @@ class hipGraphKernelNode : public hipGraphNode {
// HIP_LAUNCH_PARAM_BUFFER_SIZE, &kernargs_size,
// HIP_LAUNCH_PARAM_END }
unsigned int numExtra = 5;
pKernelParams_->extra = (void**) malloc(numExtra * sizeof(void *));
pKernelParams_->extra = (void**)malloc(numExtra * sizeof(void*));
if (pKernelParams_->extra == nullptr) {
return hipErrorOutOfMemory;
}
pKernelParams_->extra[0] = pNodeParams->extra[0];
size_t kernargs_size = *((size_t *)pNodeParams->extra[3]);
size_t kernargs_size = *((size_t*)pNodeParams->extra[3]);
pKernelParams_->extra[1] = malloc(kernargs_size);
if (pKernelParams_->extra[1] == nullptr) {
return hipErrorOutOfMemory;
}
pKernelParams_->extra[2] = pNodeParams->extra[2];
pKernelParams_->extra[3] = malloc(sizeof(void *));
pKernelParams_->extra[3] = malloc(sizeof(void*));
if (pKernelParams_->extra[3] == nullptr) {
return hipErrorOutOfMemory;
}
*((size_t *)pKernelParams_->extra[3]) = kernargs_size;
*((size_t*)pKernelParams_->extra[3]) = kernargs_size;
::memcpy(pKernelParams_->extra[1], (pNodeParams->extra[1]), kernargs_size);
pKernelParams_->extra[4] = pNodeParams->extra[4];
}
@@ -773,9 +772,7 @@ class hipGraphKernelNode : public hipGraphNode {
memset(&kernelAttr_, 0, sizeof(kernelAttr_));
}
~hipGraphKernelNode() {
freeParams();
}
~hipGraphKernelNode() { freeParams(); }
void freeParams() {
// Deallocate memory allocated for kernargs passed via 'kernelParams'
@@ -793,7 +790,7 @@ class hipGraphKernelNode : public hipGraphNode {
else {
free(pKernelParams_->extra[1]);
free(pKernelParams_->extra[3]);
memset(pKernelParams_->extra, 0, 5 * sizeof(pKernelParams_->extra[0])); // 5 items
memset(pKernelParams_->extra, 0, 5 * sizeof(pKernelParams_->extra[0])); // 5 items
free(pKernelParams_->extra);
pKernelParams_->extra = nullptr;
}
@@ -805,8 +802,7 @@ class hipGraphKernelNode : public hipGraphNode {
pKernelParams_ = new hipKernelNodeParams(*rhs.pKernelParams_);
hipError_t status = copyParams(rhs.pKernelParams_);
if (status != hipSuccess) {
ClPrint(amd::LOG_ERROR, amd::LOG_CODE,
"[hipGraph] Failed to allocate memory to copy params");
ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to allocate memory to copy params");
}
status = CopyAttr(&rhs);
if (status != hipSuccess) {
@@ -851,9 +847,9 @@ class hipGraphKernelNode : public hipGraphNode {
ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to validateKernelParams");
return status;
}
if (pKernelParams_ && ((pKernelParams_->kernelParams
&& pKernelParams_->kernelParams == params->kernelParams)
|| (pKernelParams_->extra && pKernelParams_->extra == params->extra))) {
if (pKernelParams_ &&
((pKernelParams_->kernelParams && pKernelParams_->kernelParams == params->kernelParams) ||
(pKernelParams_->extra && pKernelParams_->extra == params->extra))) {
// params is copied from pKernelParams_ and then updated, so just copy it back
*pKernelParams_ = *params;
return status;
@@ -905,7 +901,7 @@ class hipGraphKernelNode : public hipGraphNode {
return hipSuccess;
}
hipError_t CopyAttr(const hipGraphKernelNode* srcNode) {
if(srcNode->kernelAttrInUse_ != kernelAttrInUse_) {
if (srcNode->kernelAttrInUse_ != kernelAttrInUse_) {
return hipErrorInvalidContext;
}
switch (kernelAttrInUse_) {
@@ -914,7 +910,8 @@ class hipGraphKernelNode : public hipGraphNode {
kernelAttr_.accessPolicyWindow.hitProp = srcNode->kernelAttr_.accessPolicyWindow.hitProp;
kernelAttr_.accessPolicyWindow.hitRatio = srcNode->kernelAttr_.accessPolicyWindow.hitRatio;
kernelAttr_.accessPolicyWindow.missProp = srcNode->kernelAttr_.accessPolicyWindow.missProp;
kernelAttr_.accessPolicyWindow.num_bytes = srcNode->kernelAttr_.accessPolicyWindow.num_bytes;
kernelAttr_.accessPolicyWindow.num_bytes =
srcNode->kernelAttr_.accessPolicyWindow.num_bytes;
break;
case hipKernelNodeAttributeCooperative:
kernelAttr_.cooperative = srcNode->kernelAttr_.cooperative;
@@ -947,7 +944,7 @@ class hipGraphKernelNode : public hipGraphNode {
}
static hipError_t validateKernelParams(const hipKernelNodeParams* pNodeParams,
hipFunction_t *ptrFunc = nullptr, int devId = -1) {
hipFunction_t* ptrFunc = nullptr, int devId = -1) {
devId = devId == -1 ? ihipGetDevice() : devId;
hipFunction_t func = getFunc(*pNodeParams, devId);
if (!func) {
@@ -967,10 +964,9 @@ class hipGraphKernelNode : public hipGraphNode {
return status;
}
if (ptrFunc) *ptrFunc = func;
if (ptrFunc) *ptrFunc = func;
return hipSuccess;
}
};
class hipGraphMemcpyNode : public hipGraphNode {
@@ -1116,8 +1112,7 @@ class hipGraphMemcpyNode1D : public hipGraphNode {
dst_(dst),
src_(src),
count_(count),
kind_(kind) {
}
kind_(kind) {}
~hipGraphMemcpyNode1D() {}
@@ -1460,8 +1455,9 @@ class hipGraphMemsetNode : public hipGraphNode {
sprintf(buffer,
"{\n%s\n| {{ID | node handle | dptr | pitch | value | elementSize | width | "
"height} | {%u | %p | %p | %zu | %u | %u | %zu | %zu}}}",
label_.c_str(), GetID(), this, pMemsetParams_->dst, pMemsetParams_->pitch, pMemsetParams_->value,
pMemsetParams_->elementSize, pMemsetParams_->width, pMemsetParams_->height);
label_.c_str(), GetID(), this, pMemsetParams_->dst, pMemsetParams_->pitch,
pMemsetParams_->value, pMemsetParams_->elementSize, pMemsetParams_->width,
pMemsetParams_->height);
label = buffer;
} else {
size_t sizeBytes;
@@ -1470,8 +1466,8 @@ class hipGraphMemsetNode : public hipGraphNode {
} else {
sizeBytes = pMemsetParams_->width * pMemsetParams_->height;
}
label = std::to_string(GetID()) + "\n"+label_+"\n(" + std::to_string(pMemsetParams_->value) +
"," + std::to_string(sizeBytes) + ")";
label = std::to_string(GetID()) + "\n" + label_ + "\n(" +
std::to_string(pMemsetParams_->value) + "," + std::to_string(sizeBytes) + ")";
}
return label;
}
@@ -1541,7 +1537,8 @@ class hipGraphEventRecordNode : public hipGraphNode {
public:
hipGraphEventRecordNode(hipEvent_t event)
: hipGraphNode(hipGraphNodeTypeEventRecord, "solid", "rectangle", "EVENT_RECORD"), event_(event) {}
: hipGraphNode(hipGraphNodeTypeEventRecord, "solid", "rectangle", "EVENT_RECORD"),
event_(event) {}
~hipGraphEventRecordNode() {}
hipGraphNode* clone() const {
@@ -1603,7 +1600,8 @@ class hipGraphEventWaitNode : public hipGraphNode {
public:
hipGraphEventWaitNode(hipEvent_t event)
: hipGraphNode(hipGraphNodeTypeWaitEvent, "solid", "rectangle", "EVENT_WAIT"), event_(event) {}
: hipGraphNode(hipGraphNodeTypeWaitEvent, "solid", "rectangle", "EVENT_WAIT"),
event_(event) {}
~hipGraphEventWaitNode() {}
hipGraphNode* clone() const {