SWDEV-318237 SWDEV-330173 - Fixes kernel parameter corruption when multiple nodes of the same kernelNode exist

Change-Id: I6437593be46b7b5afc3d77cdd887ad4139978ad3


[ROCm/clr commit: 5549c8cd1b]

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<size_t>(pNodeParams->gridDim.x) * pNodeParams->blockDim.x;
   size_t globalWorkSizeY = static_cast<size_t>(pNodeParams->gridDim.y) * pNodeParams->blockDim.y;
   size_t globalWorkSizeZ = static_cast<size_t>(pNodeParams->gridDim.z) * pNodeParams->blockDim.z;
-  if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
-      globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
-      globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
-    return hipErrorInvalidConfiguration;
-  }
+
   status = ihipLaunchKernel_validate(
       func, static_cast<uint32_t>(globalWorkSizeX), static_cast<uint32_t>(globalWorkSizeY),
       static_cast<uint32_t>(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<size_t>(pNodeParams->gridDim.x) * pNodeParams->blockDim.x;
+  size_t globalWorkSizeY = static_cast<size_t>(pNodeParams->gridDim.y) * pNodeParams->blockDim.y;
+  size_t globalWorkSizeZ = static_cast<size_t>(pNodeParams->gridDim.z) * pNodeParams->blockDim.z;
+  if (globalWorkSizeX > std::numeric_limits<uint32_t>::max() ||
+      globalWorkSizeY > std::numeric_limits<uint32_t>::max() ||
+      globalWorkSizeZ > std::numeric_limits<uint32_t>::max()) {
+    return hipErrorInvalidConfiguration;
+  }
+
   *pGraphNode = new hipGraphKernelNode(pNodeParams, func);
   status = ihipGraphAddNode(*pGraphNode, graph, pDependencies, numDependencies);
   return status;

projects/clr/hipamd/src/hip_graph_internal.cpp

@@ -645,7 +645,7 @@ ihipGraph* ihipGraph::clone(std::unordered_map<Node, Node>& clonedNodes) const{
     newGraph->vertices_.push_back(node);
     clonedNodes[entry] = node;
   }
-  std::vector<Node> dependancies;
+
   std::vector<Node> clonedEdges;
   std::vector<Node> clonedDependencies;
   for (auto node : vertices_) {

projects/clr/hipamd/src/hip_graph_internal.hpp

@@ -236,7 +236,7 @@ struct ihipGraph {
   std::vector<Node> vertices_;
   const ihipGraph* pOriginalGraph_ = nullptr;
   static std::unordered_set<ihipGraph*> graphSet_;
-  static amd::Monitor graphSetLock_ ;
+  static amd::Monitor graphSetLock_;
 
  public:
   ihipGraph() {
@@ -295,7 +295,7 @@ struct hipGraphExec {
   std::unordered_map<Node, Node> clonedNodes_;
   amd::Command* lastEnqueuedCommand_;
   static std::unordered_set<hipGraphExec*> graphExecSet_;
-  static amd::Monitor graphExecSetLock_ ;
+  static amd::Monitor graphExecSetLock_;
 
  public:
   hipGraphExec(std::vector<Node>& levelOrder, std::vector<std::vector<Node>>& 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<hipFunction_t>(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<hipGraphKernelNode const&>(*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) {

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_) {