diff --git a/projects/clr/hipamd/src/hip_graph.cpp b/projects/clr/hipamd/src/hip_graph.cpp index c0ee11fc01..132b7d8817 100644 --- a/projects/clr/hipamd/src/hip_graph.cpp +++ b/projects/clr/hipamd/src/hip_graph.cpp @@ -43,22 +43,16 @@ inline hipError_t ihipGraphAddNode(hipGraphNode_t graphNode, hipGraph_t graph, hipError_t ihipValidateKernelParams(const hipKernelNodeParams* pNodeParams) { - if (pNodeParams->kernelParams == nullptr) { - return hipErrorInvalidValue; - } hipFunction_t func = nullptr; hipError_t status = hipGraphKernelNode::getFunc(&func, *pNodeParams, ihipGetDevice()); if (status != hipSuccess) { return hipErrorInvalidDeviceFunction; } + size_t globalWorkSizeX = static_cast(pNodeParams->gridDim.x) * pNodeParams->blockDim.x; size_t globalWorkSizeY = static_cast(pNodeParams->gridDim.y) * pNodeParams->blockDim.y; size_t globalWorkSizeZ = static_cast(pNodeParams->gridDim.z) * pNodeParams->blockDim.z; - if (globalWorkSizeX > std::numeric_limits::max() || - globalWorkSizeY > std::numeric_limits::max() || - globalWorkSizeZ > std::numeric_limits::max()) { - return hipErrorInvalidConfiguration; - } + status = ihipLaunchKernel_validate( func, static_cast(globalWorkSizeX), static_cast(globalWorkSizeY), static_cast(globalWorkSizeZ), pNodeParams->blockDim.x, pNodeParams->blockDim.y, @@ -81,15 +75,28 @@ hipError_t ihipGraphAddKernelNode(hipGraphNode_t* pGraphNode, hipGraph_t graph, if (!ihipGraph::isGraphValid(graph)) { return hipErrorInvalidValue; } - hipError_t status = ihipValidateKernelParams(pNodeParams); - if (hipSuccess != status) { - return status; - } + hipFunction_t func = nullptr; - status = hipGraphKernelNode::getFunc(&func, *pNodeParams, ihipGetDevice()); + hipError_t status = hipGraphKernelNode::getFunc(&func, *pNodeParams, ihipGetDevice()); if (status != hipSuccess) { return hipErrorInvalidDeviceFunction; } + + // If neither 'kernelParams' or 'extra' are provided or if both are provided, return error + if ((pNodeParams->kernelParams == nullptr && pNodeParams->extra == nullptr) || + (pNodeParams->kernelParams != nullptr) && (pNodeParams->extra != nullptr)) { + return hipErrorInvalidValue; + } + + size_t globalWorkSizeX = static_cast(pNodeParams->gridDim.x) * pNodeParams->blockDim.x; + size_t globalWorkSizeY = static_cast(pNodeParams->gridDim.y) * pNodeParams->blockDim.y; + size_t globalWorkSizeZ = static_cast(pNodeParams->gridDim.z) * pNodeParams->blockDim.z; + if (globalWorkSizeX > std::numeric_limits::max() || + globalWorkSizeY > std::numeric_limits::max() || + globalWorkSizeZ > std::numeric_limits::max()) { + return hipErrorInvalidConfiguration; + } + *pGraphNode = new hipGraphKernelNode(pNodeParams, func); status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies); return status; diff --git a/projects/clr/hipamd/src/hip_graph_internal.cpp b/projects/clr/hipamd/src/hip_graph_internal.cpp index dc753dcf72..76749b119c 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.cpp +++ b/projects/clr/hipamd/src/hip_graph_internal.cpp @@ -645,7 +645,7 @@ ihipGraph* ihipGraph::clone(std::unordered_map& clonedNodes) const{ newGraph->vertices_.push_back(node); clonedNodes[entry] = node; } - std::vector dependancies; + std::vector clonedEdges; std::vector clonedDependencies; for (auto node : vertices_) { diff --git a/projects/clr/hipamd/src/hip_graph_internal.hpp b/projects/clr/hipamd/src/hip_graph_internal.hpp index 3e3f66165b..65197b0b94 100644 --- a/projects/clr/hipamd/src/hip_graph_internal.hpp +++ b/projects/clr/hipamd/src/hip_graph_internal.hpp @@ -236,7 +236,7 @@ struct ihipGraph { std::vector vertices_; const ihipGraph* pOriginalGraph_ = nullptr; static std::unordered_set graphSet_; - static amd::Monitor graphSetLock_ ; + static amd::Monitor graphSetLock_; public: ihipGraph() { @@ -295,7 +295,7 @@ struct hipGraphExec { std::unordered_map clonedNodes_; amd::Command* lastEnqueuedCommand_; static std::unordered_set graphExecSet_; - static amd::Monitor graphExecSetLock_ ; + static amd::Monitor graphExecSetLock_; public: hipGraphExec(std::vector& levelOrder, std::vector>& lists, @@ -309,7 +309,7 @@ struct hipGraphExec { currentQueueIndex_(0) { amd::ScopedLock lock(graphExecSetLock_); graphExecSet_.insert(this); - } + } ~hipGraphExec() { // new commands are launched for every launch they are destroyed as and when command is @@ -456,8 +456,11 @@ struct hipChildGraphNode : public hipGraphNode { class hipGraphKernelNode : public hipGraphNode { hipKernelNodeParams* pKernelParams_; hipFunction_t func_; + unsigned int numParams_; + public: - static hipError_t getFunc(hipFunction_t* func, const hipKernelNodeParams& params, unsigned int device) { + static hipError_t getFunc(hipFunction_t* func, const hipKernelNodeParams& params, + unsigned int device) { hipError_t status = PlatformState::instance().getStatFunc(func, params.func, device); if (status != hipSuccess) { *func = reinterpret_cast(params.func); @@ -467,21 +470,109 @@ class hipGraphKernelNode : public hipGraphNode { } return hipSuccess; } + + hipError_t copyParams(const hipKernelNodeParams* pNodeParams, const hipFunction_t 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*)); + 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_); + } + } + + // Allocate/assign memory if params are passed as part of 'extra' + else if (pNodeParams->extra != nullptr) { + // 'extra' is a struct that contains the following info: { + // HIP_LAUNCH_PARAM_BUFFER_POINTER, kernargs, + // 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) { + return hipErrorOutOfMemory; + } + pKernelParams_->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) { + return hipErrorOutOfMemory; + } + pKernelParams_->extra[2] = pNodeParams->extra[2]; + pKernelParams_->extra[3] = malloc(sizeof(void *)); + if (pKernelParams_->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]; + } + return hipSuccess; + } + hipGraphKernelNode(const hipKernelNodeParams* pNodeParams, const hipFunction_t func) : hipGraphNode(hipGraphNodeTypeKernel) { pKernelParams_ = new hipKernelNodeParams(*pNodeParams); func_ = func; + hipError_t status = copyParams(pNodeParams, func_); + if (status != hipSuccess) { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, + "[hipGraph] Failed to allocate memory to copy kernel arguments"); + } } - ~hipGraphKernelNode() { delete pKernelParams_; } + ~hipGraphKernelNode() { + // Deallocate memory allocated for kernargs passed via 'kernelParams' + if (pKernelParams_->kernelParams != nullptr) { + for (size_t i = 0; i < numParams_; ++i) { + if (pKernelParams_->kernelParams[i] != nullptr) { + free(pKernelParams_->kernelParams[i]); + } + pKernelParams_->kernelParams[i] = nullptr; + } + free(pKernelParams_->kernelParams); + pKernelParams_->kernelParams = nullptr; + } + // Deallocate memory allocated for kernargs passed via 'extra' + else { + free(pKernelParams_->extra[1]); + free(pKernelParams_->extra[3]); + free(pKernelParams_->extra); + } + delete pKernelParams_; + } + hipGraphKernelNode(const hipGraphKernelNode& rhs) : hipGraphNode(rhs) { pKernelParams_ = new hipKernelNodeParams(*rhs.pKernelParams_); func_ = rhs.func_; + numParams_ = rhs.numParams_; + hipError_t status = copyParams(rhs.pKernelParams_, func_); + if (status != hipSuccess) { + ClPrint(amd::LOG_ERROR, amd::LOG_CODE, + "[hipGraph] Failed to allocate memory to deep copy kernargs"); + } } hipGraphNode* clone() const { return new hipGraphKernelNode(static_cast(*this)); } hipError_t CreateCommand(amd::HostQueue* queue) { - hipError_t status = hipGraphNode::CreateCommand(queue); + hipError_t status = ihipValidateKernelParams(pKernelParams_); + if (hipSuccess != status) { + return status; + } + status = hipGraphNode::CreateCommand(queue); if (status != hipSuccess) { return status; } @@ -506,6 +597,11 @@ class hipGraphKernelNode : public hipGraphNode { if (hipSuccess != status) { return status; } + + hip::DeviceFunc* oldFunction = hip::DeviceFunc::asFunction(func_); + amd::Kernel* oldKernel = oldFunction->kernel(); + const amd::KernelSignature& oldSignature = oldKernel->signature(); + if (params->func != pKernelParams_->func) { hipFunction_t func = nullptr; hipError_t status = hipGraphKernelNode::getFunc(&func, *params, ihipGetDevice()); @@ -514,7 +610,81 @@ class hipGraphKernelNode : public hipGraphNode { } func_ = func; } - std::memcpy(pKernelParams_, params, sizeof(hipKernelNodeParams)); + + hip::DeviceFunc* newFunction = hip::DeviceFunc::asFunction(func_); + amd::Kernel* newKernel = newFunction->kernel(); + const amd::KernelSignature& newSignature = newKernel->signature(); + unsigned int numParams = newSignature.numParameters(); + // When params are passed in 'kernelParams' reallocate for updated func and its params + if (params->kernelParams != nullptr) { + if (numParams != numParams_) { + pKernelParams_->kernelParams = (void **) realloc(pKernelParams_->kernelParams, + numParams * sizeof(void *)); + if (pKernelParams_->kernelParams == nullptr) { + return hipErrorOutOfMemory; + } + } + + uint32_t i = 0; + int minParams = std::min(numParams, numParams_); + // Rewrite the contents in kernelParams + for (i = 0; i < minParams; ++i) { + const amd::KernelParameterDescriptor& newDesc = newSignature.at(i); + const amd::KernelParameterDescriptor& oldDesc = oldSignature.at(i); + + if (newDesc.size_ != oldDesc.size_) { + pKernelParams_->kernelParams[i] = (void *) realloc(pKernelParams_->kernelParams[i], + newDesc.size_); + if (pKernelParams_->kernelParams[i] == nullptr) { + return hipErrorOutOfMemory; + } + } + ::memcpy(pKernelParams_->kernelParams[i], (params->kernelParams[i]), newDesc.size_); + } + // If the numParams of the new func is greater than old func, allocate + if ((numParams > numParams_)) { + const amd::KernelParameterDescriptor& newDesc = newSignature.at(i); + for (;i < numParams; i++) { + pKernelParams_->kernelParams[i] = (void *) malloc(newDesc.size_); + if (pKernelParams_->kernelParams[i] == nullptr) { + return hipErrorOutOfMemory; + } + ::memcpy(pKernelParams_->kernelParams[i], (params->kernelParams[i]), newDesc.size_); + } + } + // If the numParams of old func is greater than new func, free + else if ((numParams < numParams_) && (numParams != numParams_)) { + for (; i < numParams_; i++) { + if (pKernelParams_->kernelParams[i] != nullptr) { + free(pKernelParams_->kernelParams[i]); + } + pKernelParams_->kernelParams[i] = nullptr; + } + } + pKernelParams_->extra = params->extra; + } + // When params are passed as a part of 'extra', reallocate for updated func and its params + else if (params->extra != nullptr) { + + pKernelParams_->extra[0] = params->extra[0]; + size_t kernargs_size_new = *((size_t *)params->extra[3]); + size_t kernargs_size_old = *((size_t *)pKernelParams_->extra[3]); + if (kernargs_size_new != kernargs_size_old) { + pKernelParams_->extra[1] = (void *) realloc(pKernelParams_->extra[1], kernargs_size_new); + if (pKernelParams_->extra[1] == nullptr) { + return hipErrorOutOfMemory; + } + } + pKernelParams_->extra[2] = params->extra[2]; + *((size_t *)pKernelParams_->extra[3]) = kernargs_size_new; + ::memcpy(pKernelParams_->extra[1], (params->extra[1]), kernargs_size_new); + pKernelParams_->extra[4] = params->extra[4]; + } + numParams_ = numParams; + pKernelParams_->blockDim = params->blockDim; + pKernelParams_->func = params->func; + pKernelParams_->gridDim = params->gridDim; + pKernelParams_->sharedMemBytes = params->sharedMemBytes; return status; } hipError_t SetAttrParams(hipKernelNodeAttrID attr, const hipKernelNodeAttrValue* params) { @@ -735,7 +905,7 @@ class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D { hipMemcpyKind kind) { size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; - //check to see if dst is also a symbol (cuda negative test case) + // check to see if dst is also a symbol (cuda negative test case) hipError_t status = ihipMemcpySymbol_validate(dst, count, offset, sym_size, device_ptr); if (status == hipSuccess) { return hipErrorInvalidValue; @@ -747,11 +917,11 @@ class hipGraphMemcpyNodeFromSymbol : public hipGraphMemcpyNode1D { size_t dOffset = 0; amd::Memory* dstMemory = getMemoryObject(dst, dOffset); - if( dstMemory == nullptr && kind != hipMemcpyHostToDevice) { + if (dstMemory == nullptr && kind != hipMemcpyHostToDevice) { return hipErrorInvalidMemcpyDirection; - } else if ( dstMemory != nullptr && kind != hipMemcpyDeviceToDevice) { + } else if (dstMemory != nullptr && kind != hipMemcpyDeviceToDevice) { return hipErrorInvalidMemcpyDirection; - } else if ( kind == hipMemcpyHostToHost || kind == hipMemcpyDeviceToHost) { + } else if (kind == hipMemcpyHostToHost || kind == hipMemcpyDeviceToHost) { return hipErrorInvalidMemcpyDirection; } @@ -823,10 +993,9 @@ class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D { hipError_t SetParams(const void* symbol, const void* src, size_t count, size_t offset, hipMemcpyKind kind) { - size_t sym_size = 0; hipDeviceptr_t device_ptr = nullptr; - //check to see if src is also a symbol (cuda negative test case) + // check to see if src is also a symbol (cuda negative test case) hipError_t status = ihipMemcpySymbol_validate(src, count, offset, sym_size, device_ptr); if (status == hipSuccess) { return hipErrorInvalidValue; @@ -837,7 +1006,7 @@ class hipGraphMemcpyNodeToSymbol : public hipGraphMemcpyNode1D { } size_t dOffset = 0; amd::Memory* srcMemory = getMemoryObject(src, dOffset); - if(srcMemory == nullptr && kind != hipMemcpyHostToDevice) { + if (srcMemory == nullptr && kind != hipMemcpyHostToDevice) { return hipErrorInvalidValue; } else if (srcMemory != nullptr && kind != hipMemcpyDeviceToDevice) { return hipErrorInvalidValue; @@ -912,7 +1081,6 @@ class hipGraphMemsetNode : public hipGraphNode { std::memcpy(params, pMemsetParams_, sizeof(hipMemsetParams)); } hipError_t SetParams(const hipMemsetParams* params) { - hipError_t hip_error = hipSuccess; hip_error = ihipGraphMemsetParams_validate(params); if (hip_error != hipSuccess) { @@ -923,9 +1091,9 @@ class hipGraphMemsetNode : public hipGraphNode { params->width * params->elementSize); } else { auto sizeBytes = params->width * params->height * 1; - hip_error = ihipMemset3D_validate( - {params->dst, params->pitch, params->width, params->height}, - params->value, {params->width, params->height, 1}, sizeBytes); + hip_error = + ihipMemset3D_validate({params->dst, params->pitch, params->width, params->height}, + params->value, {params->width, params->height, 1}, sizeBytes); } if (hip_error != hipSuccess) { diff --git a/projects/clr/hipamd/src/hip_module.cpp b/projects/clr/hipamd/src/hip_module.cpp index d527d4f733..4518f4589e 100644 --- a/projects/clr/hipamd/src/hip_module.cpp +++ b/projects/clr/hipamd/src/hip_module.cpp @@ -217,13 +217,6 @@ hipError_t ihipLaunchKernel_validate(hipFunction_t f, uint32_t globalWorkSizeX, return hipErrorInvalidValue; } - if (extra != nullptr) { - if (extra[0] != HIP_LAUNCH_PARAM_BUFFER_POINTER || extra[2] != HIP_LAUNCH_PARAM_BUFFER_SIZE || - extra[4] != HIP_LAUNCH_PARAM_END) { - return hipErrorNotInitialized; - } - } - const amd::Device* device = g_devices[deviceId]->devices()[0]; // Make sure dispatch doesn't exceed max workgroup size limit if (blockDimX * blockDimY * blockDimZ > device->info().maxWorkGroupSize_) {