diff --git a/hipamd/src/hip_graph_internal.hpp b/hipamd/src/hip_graph_internal.hpp index 43dd4e4202..6ee9f315e1 100644 --- a/hipamd/src/hip_graph_internal.hpp +++ b/hipamd/src/hip_graph_internal.hpp @@ -731,7 +731,7 @@ struct hipChildGraphNode : public hipGraphNode { }; class hipGraphKernelNode : public hipGraphNode { - hipKernelNodeParams* pKernelParams_; + hipKernelNodeParams kernelParams_; unsigned int numParams_; hipKernelNodeAttrValue kernelAttr_; unsigned int kernelAttrInUse_; @@ -757,7 +757,7 @@ class hipGraphKernelNode : public hipGraphNode { } std::string GetLabel(hipGraphDebugDotFlags flag) { - hipFunction_t func = getFunc(*pKernelParams_, ihipGetDevice()); + hipFunction_t func = getFunc(kernelParams_, ihipGetDevice()); hip::DeviceFunc* function = hip::DeviceFunc::asFunction(func); std::string label; char buffer[500]; @@ -767,10 +767,10 @@ class hipGraphKernelNode : public hipGraphNode { "handle | func handle} | {%p | %p}}\n| {accessPolicyWindow | {base_ptr | num_bytes | " "hitRatio | hitProp | missProp} | {%p | %zu | %f | %d | %d}}\n| {cooperative | " "%u}\n| {priority | 0}\n}", - label_.c_str(), 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, this, pKernelParams_->func, + label_.c_str(), GetID(), function->name().c_str(), kernelParams_.gridDim.x, + kernelParams_.gridDim.y, kernelParams_.gridDim.z, kernelParams_.blockDim.x, + kernelParams_.blockDim.y, kernelParams_.blockDim.z, + kernelParams_.sharedMemBytes, this, kernelParams_.func, kernelAttr_.accessPolicyWindow.base_ptr, kernelAttr_.accessPolicyWindow.num_bytes, kernelAttr_.accessPolicyWindow.hitRatio, kernelAttr_.accessPolicyWindow.hitProp, kernelAttr_.accessPolicyWindow.missProp, kernelAttr_.cooperative); @@ -790,10 +790,10 @@ class hipGraphKernelNode : public hipGraphNode { } 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); + GetID(), function->name().c_str(), kernelParams_.gridDim.x, + kernelParams_.gridDim.y, kernelParams_.gridDim.z, + kernelParams_.blockDim.x, kernelParams_.blockDim.y, + kernelParams_.blockDim.z, kernelParams_.sharedMemBytes); label = buffer; } else { @@ -837,18 +837,18 @@ class hipGraphKernelNode : public hipGraphNode { // Allocate/assign memory if params are passed part of 'kernelParams' if (pNodeParams->kernelParams != nullptr) { - pKernelParams_->kernelParams = (void**)malloc(numParams_ * sizeof(void*)); - if (pKernelParams_->kernelParams == nullptr) { + kernelParams_.kernelParams = (void**)malloc(numParams_ * sizeof(void*)); + if (kernelParams_.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) { + kernelParams_.kernelParams[i] = malloc(desc.size_); + if (kernelParams_.kernelParams[i] == nullptr) { return hipErrorOutOfMemory; } - ::memcpy(pKernelParams_->kernelParams[i], (pNodeParams->kernelParams[i]), desc.size_); + ::memcpy(kernelParams_.kernelParams[i], (pNodeParams->kernelParams[i]), desc.size_); } } @@ -859,31 +859,31 @@ 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*)); - if (pKernelParams_->extra == nullptr) { + kernelParams_.extra = (void**)malloc(numExtra * sizeof(void*)); + if (kernelParams_.extra == nullptr) { return hipErrorOutOfMemory; } - pKernelParams_->extra[0] = pNodeParams->extra[0]; + kernelParams_.extra[0] = pNodeParams->extra[0]; size_t kernargs_size = *((size_t*)pNodeParams->extra[3]); - pKernelParams_->extra[1] = malloc(kernargs_size); - if (pKernelParams_->extra[1] == nullptr) { + kernelParams_.extra[1] = malloc(kernargs_size); + if (kernelParams_.extra[1] == nullptr) { return hipErrorOutOfMemory; } - pKernelParams_->extra[2] = pNodeParams->extra[2]; - pKernelParams_->extra[3] = malloc(sizeof(void*)); - if (pKernelParams_->extra[3] == nullptr) { + kernelParams_.extra[2] = pNodeParams->extra[2]; + kernelParams_.extra[3] = malloc(sizeof(void*)); + if (kernelParams_.extra[3] == nullptr) { return hipErrorOutOfMemory; } - *((size_t*)pKernelParams_->extra[3]) = kernargs_size; - ::memcpy(pKernelParams_->extra[1], (pNodeParams->extra[1]), kernargs_size); - pKernelParams_->extra[4] = pNodeParams->extra[4]; + *((size_t*)kernelParams_.extra[3]) = kernargs_size; + ::memcpy(kernelParams_.extra[1], (pNodeParams->extra[1]), kernargs_size); + kernelParams_.extra[4] = pNodeParams->extra[4]; } return hipSuccess; } hipGraphKernelNode(const hipKernelNodeParams* pNodeParams) : hipGraphNode(hipGraphNodeTypeKernel, "bold", "octagon", "KERNEL") { - pKernelParams_ = new hipKernelNodeParams(*pNodeParams); + kernelParams_ = *pNodeParams; if (copyParams(pNodeParams) != hipSuccess) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to copy params"); } @@ -895,31 +895,29 @@ class hipGraphKernelNode : public hipGraphNode { void freeParams() { // Deallocate memory allocated for kernargs passed via 'kernelParams' - if (pKernelParams_->kernelParams != nullptr) { + if (kernelParams_.kernelParams != nullptr) { for (size_t i = 0; i < numParams_; ++i) { - if (pKernelParams_->kernelParams[i] != nullptr) { - free(pKernelParams_->kernelParams[i]); + if (kernelParams_.kernelParams[i] != nullptr) { + free(kernelParams_.kernelParams[i]); } - pKernelParams_->kernelParams[i] = nullptr; + kernelParams_.kernelParams[i] = nullptr; } - free(pKernelParams_->kernelParams); - pKernelParams_->kernelParams = nullptr; + free(kernelParams_.kernelParams); + kernelParams_.kernelParams = nullptr; } // Deallocate memory allocated for kernargs passed via 'extra' else { - free(pKernelParams_->extra[1]); - free(pKernelParams_->extra[3]); - memset(pKernelParams_->extra, 0, 5 * sizeof(pKernelParams_->extra[0])); // 5 items - free(pKernelParams_->extra); - pKernelParams_->extra = nullptr; + free(kernelParams_.extra[1]); + free(kernelParams_.extra[3]); + memset(kernelParams_.extra, 0, 5 * sizeof(kernelParams_.extra[0])); // 5 items + free(kernelParams_.extra); + kernelParams_.extra = nullptr; } - delete pKernelParams_; - pKernelParams_ = nullptr; } hipGraphKernelNode(const hipGraphKernelNode& rhs) : hipGraphNode(rhs) { - pKernelParams_ = new hipKernelNodeParams(*rhs.pKernelParams_); - hipError_t status = copyParams(rhs.pKernelParams_); + kernelParams_ = rhs.kernelParams_; + hipError_t status = copyParams(&rhs.kernelParams_); if (status != hipSuccess) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to allocate memory to copy params"); } @@ -937,7 +935,7 @@ class hipGraphKernelNode : public hipGraphNode { hipError_t CreateCommand(hip::Stream* stream) { hipFunction_t func = nullptr; - hipError_t status = validateKernelParams(pKernelParams_, &func, + hipError_t status = validateKernelParams(&kernelParams_, &func, stream ? hip::getDeviceID(stream->context()) : -1); if (hipSuccess != status) { return status; @@ -949,17 +947,17 @@ class hipGraphKernelNode : public hipGraphNode { commands_.reserve(1); amd::Command* command; status = ihipLaunchKernelCommand( - command, func, pKernelParams_->gridDim.x * pKernelParams_->blockDim.x, - pKernelParams_->gridDim.y * pKernelParams_->blockDim.y, - pKernelParams_->gridDim.z * pKernelParams_->blockDim.z, pKernelParams_->blockDim.x, - pKernelParams_->blockDim.y, pKernelParams_->blockDim.z, pKernelParams_->sharedMemBytes, - stream, pKernelParams_->kernelParams, pKernelParams_->extra, nullptr, nullptr, 0, 0, 0, 0, 0, + command, func, kernelParams_.gridDim.x * kernelParams_.blockDim.x, + kernelParams_.gridDim.y * kernelParams_.blockDim.y, + kernelParams_.gridDim.z * kernelParams_.blockDim.z, kernelParams_.blockDim.x, + kernelParams_.blockDim.y, kernelParams_.blockDim.z, kernelParams_.sharedMemBytes, + stream, kernelParams_.kernelParams, kernelParams_.extra, nullptr, nullptr, 0, 0, 0, 0, 0, 0, 0); commands_.emplace_back(command); return status; } - void GetParams(hipKernelNodeParams* params) { *params = *pKernelParams_; } + void GetParams(hipKernelNodeParams* params) { *params = kernelParams_; } hipError_t SetParams(const hipKernelNodeParams* params) { // updates kernel params @@ -968,15 +966,14 @@ 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))) { - // params is copied from pKernelParams_ and then updated, so just copy it back - *pKernelParams_ = *params; + if ((kernelParams_.kernelParams && kernelParams_.kernelParams == params->kernelParams) || + (kernelParams_.extra && kernelParams_.extra == params->extra)) { + // params is copied from kernelParams_ and then updated, so just copy it back + kernelParams_ = *params; return status; } freeParams(); - pKernelParams_ = new hipKernelNodeParams(*params); + kernelParams_ = *params; status = copyParams(params); if (status != hipSuccess) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed to set params"); @@ -1059,7 +1056,7 @@ class hipGraphKernelNode : public hipGraphNode { hipError_t SetParams(hipGraphNode* node) { const hipGraphKernelNode* kernelNode = static_cast(node); - return SetParams(kernelNode->pKernelParams_); + return SetParams(&kernelNode->kernelParams_); } static hipError_t validateKernelParams(const hipKernelNodeParams* pNodeParams, @@ -1089,17 +1086,17 @@ class hipGraphKernelNode : public hipGraphNode { }; class hipGraphMemcpyNode : public hipGraphNode { - hipMemcpy3DParms* pCopyParams_; + hipMemcpy3DParms copyParams_; public: hipGraphMemcpyNode(const hipMemcpy3DParms* pCopyParams) : hipGraphNode(hipGraphNodeTypeMemcpy, "solid", "trapezium", "MEMCPY") { - pCopyParams_ = new hipMemcpy3DParms(*pCopyParams); + copyParams_ = *pCopyParams; } - ~hipGraphMemcpyNode() { delete pCopyParams_; } + ~hipGraphMemcpyNode() {} hipGraphMemcpyNode(const hipGraphMemcpyNode& rhs) : hipGraphNode(rhs) { - pCopyParams_ = new hipMemcpy3DParms(*rhs.pCopyParams_); + copyParams_ = rhs.copyParams_; } hipGraphNode* clone() const { @@ -1107,7 +1104,7 @@ class hipGraphMemcpyNode : public hipGraphNode { } hipError_t CreateCommand(hip::Stream* stream) { - if (IsHtoHMemcpy(pCopyParams_->dstPtr.ptr, pCopyParams_->srcPtr.ptr, pCopyParams_->kind)) { + if (IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.kind)) { return hipSuccess; } hipError_t status = hipGraphNode::CreateCommand(stream); @@ -1116,42 +1113,42 @@ class hipGraphMemcpyNode : public hipGraphNode { } commands_.reserve(1); amd::Command* command; - status = ihipMemcpy3DCommand(command, pCopyParams_, stream); + status = ihipMemcpy3DCommand(command, ©Params_, stream); commands_.emplace_back(command); return status; } void EnqueueCommands(hipStream_t stream) override { - if (isEnabled_ && IsHtoHMemcpy(pCopyParams_->dstPtr.ptr, pCopyParams_->srcPtr.ptr, pCopyParams_->kind)) { - ihipHtoHMemcpy(pCopyParams_->dstPtr.ptr, pCopyParams_->srcPtr.ptr, - pCopyParams_->extent.width * pCopyParams_->extent.height * - pCopyParams_->extent.depth, *hip::getStream(stream)); + if (isEnabled_ && IsHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, copyParams_.kind)) { + ihipHtoHMemcpy(copyParams_.dstPtr.ptr, copyParams_.srcPtr.ptr, + copyParams_.extent.width * copyParams_.extent.height * + copyParams_.extent.depth, *hip::getStream(stream)); return; } hipGraphNode::EnqueueCommands(stream); } void GetParams(hipMemcpy3DParms* params) { - std::memcpy(params, pCopyParams_, sizeof(hipMemcpy3DParms)); + std::memcpy(params, ©Params_, sizeof(hipMemcpy3DParms)); } hipError_t SetParams(const hipMemcpy3DParms* params) { hipError_t status = ValidateParams(params); if (status != hipSuccess) { return status; } - std::memcpy(pCopyParams_, params, sizeof(hipMemcpy3DParms)); + std::memcpy(©Params_, params, sizeof(hipMemcpy3DParms)); return hipSuccess; } hipError_t SetParams(hipGraphNode* node) { const hipGraphMemcpyNode* memcpyNode = static_cast(node); - return SetParams(memcpyNode->pCopyParams_); + return SetParams(&memcpyNode->copyParams_); } // ToDo: use this when commands are cloned and command params are to be updated hipError_t ValidateParams(const hipMemcpy3DParms* pNodeParams); std::string GetLabel(hipGraphDebugDotFlags flag) { size_t offset = 0; - const HIP_MEMCPY3D pCopy = hip::getDrvMemcpy3DDesc(*pCopyParams_); + const HIP_MEMCPY3D pCopy = hip::getDrvMemcpy3DDesc(copyParams_); hipMemoryType srcMemoryType = pCopy.srcMemoryType; if (srcMemoryType == hipMemoryTypeUnified) { srcMemoryType = @@ -1211,13 +1208,13 @@ class hipGraphMemcpyNode : public hipGraphNode { "| %zu}}\n| {{srcPos | {{x | %zu} | {y | %zu} | {z | %zu}}} | {dstPos | {{x | %zu} | {y " "| " "%zu} | {z | %zu}}} | {Extent | {{Width | %zu} | {Height | %zu} | {Depth | %zu}}}}\n}", - label_.c_str(), GetID(), this, memcpyDirection.c_str(), pCopyParams_->srcPtr.pitch, - pCopyParams_->srcPtr.ptr, pCopyParams_->srcPtr.xsize, pCopyParams_->srcPtr.ysize, - pCopyParams_->dstPtr.pitch, pCopyParams_->dstPtr.ptr, pCopyParams_->dstPtr.xsize, - pCopyParams_->dstPtr.ysize, pCopyParams_->srcPos.x, pCopyParams_->srcPos.y, - pCopyParams_->srcPos.z, pCopyParams_->dstPos.x, pCopyParams_->dstPos.y, - pCopyParams_->dstPos.z, pCopyParams_->extent.width, pCopyParams_->extent.height, - pCopyParams_->extent.depth); + label_.c_str(), GetID(), this, memcpyDirection.c_str(), copyParams_.srcPtr.pitch, + copyParams_.srcPtr.ptr, copyParams_.srcPtr.xsize, copyParams_.srcPtr.ysize, + copyParams_.dstPtr.pitch, copyParams_.dstPtr.ptr, copyParams_.dstPtr.xsize, + copyParams_.dstPtr.ysize, copyParams_.srcPos.x, copyParams_.srcPos.y, + copyParams_.srcPos.z, copyParams_.dstPos.x, copyParams_.dstPos.y, + copyParams_.dstPos.z, copyParams_.extent.width, copyParams_.extent.height, + copyParams_.extent.depth); label = buffer; } else { label = std::to_string(GetID()) + "\nMEMCPY\n(" + memcpyDirection + ")"; @@ -1582,24 +1579,24 @@ class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D { }; class hipGraphMemsetNode : public hipGraphNode { - hipMemsetParams* pMemsetParams_; + hipMemsetParams memsetParams_; public: hipGraphMemsetNode(const hipMemsetParams* pMemsetParams) : hipGraphNode(hipGraphNodeTypeMemset, "solid", "invtrapezium", "MEMSET") { - pMemsetParams_ = new hipMemsetParams(*pMemsetParams); + memsetParams_ = *pMemsetParams; size_t sizeBytes = 0; - if (pMemsetParams_->height == 1) { - sizeBytes = pMemsetParams_->width * pMemsetParams_->elementSize; + if (memsetParams_.height == 1) { + sizeBytes = memsetParams_.width * memsetParams_.elementSize; } else { - sizeBytes = pMemsetParams_->width * pMemsetParams_->height * pMemsetParams_->elementSize; + sizeBytes = memsetParams_.width * memsetParams_.height * memsetParams_.elementSize; } } - ~hipGraphMemsetNode() { delete pMemsetParams_; } + ~hipGraphMemsetNode() { } // Copy constructor hipGraphMemsetNode(const hipGraphMemsetNode& memsetNode) : hipGraphNode(memsetNode) { - pMemsetParams_ = new hipMemsetParams(*memsetNode.pMemsetParams_); + memsetParams_ = memsetNode.memsetParams_; } hipGraphNode* clone() const { @@ -1613,19 +1610,19 @@ 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, memsetParams_.dst, memsetParams_.pitch, + memsetParams_.value, memsetParams_.elementSize, memsetParams_.width, + memsetParams_.height); label = buffer; } else { size_t sizeBytes; - if (pMemsetParams_->height == 1) { - sizeBytes = pMemsetParams_->width * pMemsetParams_->elementSize; + if (memsetParams_.height == 1) { + sizeBytes = memsetParams_.width * memsetParams_.elementSize; } else { - sizeBytes = pMemsetParams_->width * pMemsetParams_->height * pMemsetParams_->elementSize; + sizeBytes = memsetParams_.width * memsetParams_.height * memsetParams_.elementSize; } label = std::to_string(GetID()) + "\n" + label_ + "\n(" + - std::to_string(pMemsetParams_->value) + "," + std::to_string(sizeBytes) + ")"; + std::to_string(memsetParams_.value) + "," + std::to_string(sizeBytes) + ")"; } return label; } @@ -1643,22 +1640,22 @@ class hipGraphMemsetNode : public hipGraphNode { if (status != hipSuccess) { return status; } - if (pMemsetParams_->height == 1) { - size_t sizeBytes = pMemsetParams_->width * pMemsetParams_->elementSize; - hipError_t status = ihipMemsetCommand(commands_, pMemsetParams_->dst, pMemsetParams_->value, - pMemsetParams_->elementSize, sizeBytes, stream); + if (memsetParams_.height == 1) { + size_t sizeBytes = memsetParams_.width * memsetParams_.elementSize; + hipError_t status = ihipMemsetCommand(commands_, memsetParams_.dst, memsetParams_.value, + memsetParams_.elementSize, sizeBytes, stream); } else { hipError_t status = ihipMemset3DCommand( commands_, - {pMemsetParams_->dst, pMemsetParams_->pitch, pMemsetParams_->width * pMemsetParams_->elementSize, - pMemsetParams_->height}, - pMemsetParams_->value, {pMemsetParams_->width * pMemsetParams_->elementSize, pMemsetParams_->height, 1}, stream, pMemsetParams_->elementSize); + {memsetParams_.dst, memsetParams_.pitch, memsetParams_.width * memsetParams_.elementSize, + memsetParams_.height}, + memsetParams_.value, {memsetParams_.width * memsetParams_.elementSize, memsetParams_.height, 1}, stream, memsetParams_.elementSize); } return status; } void GetParams(hipMemsetParams* params) { - std::memcpy(params, pMemsetParams_, sizeof(hipMemsetParams)); + std::memcpy(params, &memsetParams_, sizeof(hipMemsetParams)); } hipError_t SetParams(const hipMemsetParams* params, bool isExec = false) { @@ -1671,7 +1668,7 @@ class hipGraphMemsetNode : public hipGraphNode { size_t discardOffset = 0; amd::Memory *memObj = getMemoryObject(params->dst, discardOffset); if (memObj != nullptr) { - amd::Memory *memObjOri = getMemoryObject(pMemsetParams_->dst, discardOffset); + amd::Memory *memObjOri = getMemoryObject(memsetParams_.dst, discardOffset); if (memObjOri != nullptr) { if (memObjOri->getUserData().deviceId != memObj->getUserData().deviceId) { return hipErrorInvalidValue; @@ -1696,8 +1693,8 @@ class hipGraphMemsetNode : public hipGraphNode { if (isExec) { // 2D - hipGraphExecMemsetNodeSetParams returns invalid value if new width or new height is // not same as what memset node is added with. - if (pMemsetParams_->width * pMemsetParams_->elementSize != params->width * params->elementSize - || pMemsetParams_->height != params->height) { + if (memsetParams_.width * memsetParams_.elementSize != params->width * params->elementSize + || memsetParams_.height != params->height) { return hipErrorInvalidValue; } } else { @@ -1720,13 +1717,13 @@ class hipGraphMemsetNode : public hipGraphNode { if (hip_error != hipSuccess) { return hip_error; } - std::memcpy(pMemsetParams_, params, sizeof(hipMemsetParams)); + std::memcpy(&memsetParams_, params, sizeof(hipMemsetParams)); return hipSuccess; } hipError_t SetParams(hipGraphNode* node) { const hipGraphMemsetNode* memsetNode = static_cast(node); - return SetParams(memsetNode->pMemsetParams_); + return SetParams(&memsetNode->memsetParams_); } }; @@ -1836,17 +1833,17 @@ class hipGraphEventWaitNode : public hipGraphNode { }; class hipGraphHostNode : public hipGraphNode { - hipHostNodeParams* pNodeParams_; + hipHostNodeParams NodeParams_; public: - hipGraphHostNode(const hipHostNodeParams* pNodeParams) + hipGraphHostNode(const hipHostNodeParams* NodeParams) : hipGraphNode(hipGraphNodeTypeHost, "solid", "rectangle", "HOST") { - pNodeParams_ = new hipHostNodeParams(*pNodeParams); + NodeParams_ = *NodeParams; } - ~hipGraphHostNode() { delete pNodeParams_; } + ~hipGraphHostNode() { } hipGraphHostNode(const hipGraphHostNode& hostNode) : hipGraphNode(hostNode) { - pNodeParams_ = new hipHostNodeParams(*hostNode.pNodeParams_); + NodeParams_ = hostNode.NodeParams_; } hipGraphNode* clone() const { @@ -1866,13 +1863,13 @@ class hipGraphHostNode : public hipGraphNode { } static void Callback(cl_event event, cl_int command_exec_status, void* user_data) { - hipHostNodeParams* pNodeParams = reinterpret_cast(user_data); - pNodeParams->fn(pNodeParams->userData); + hipHostNodeParams* NodeParams = reinterpret_cast(user_data); + NodeParams->fn(NodeParams->userData); } void EnqueueCommands(hipStream_t stream) { if (!commands_.empty()) { - if (!commands_[0]->setCallback(CL_COMPLETE, hipGraphHostNode::Callback, pNodeParams_)) { + if (!commands_[0]->setCallback(CL_COMPLETE, hipGraphHostNode::Callback, &NodeParams_)) { ClPrint(amd::LOG_ERROR, amd::LOG_CODE, "[hipGraph] Failed during setCallback"); } commands_[0]->enqueue(); @@ -1891,16 +1888,16 @@ class hipGraphHostNode : public hipGraphNode { } void GetParams(hipHostNodeParams* params) { - std::memcpy(params, pNodeParams_, sizeof(hipHostNodeParams)); + std::memcpy(params, &NodeParams_, sizeof(hipHostNodeParams)); } hipError_t SetParams(const hipHostNodeParams* params) { - std::memcpy(pNodeParams_, params, sizeof(hipHostNodeParams)); + std::memcpy(&NodeParams_, params, sizeof(hipHostNodeParams)); return hipSuccess; } hipError_t SetParams(hipGraphNode* node) { const hipGraphHostNode* hostNode = static_cast(node); - return SetParams(hostNode->pNodeParams_); + return SetParams(&hostNode->NodeParams_); } };