SWDEV-331066 - support for LimitStackSize

Change-Id: Ie6ae74f008b4f72de83663194aafb0ebdddfc8b6


[ROCm/clr commit: 51a00aeefe]

projects/clr/rocclr/device/device.cpp

@@ -718,6 +718,15 @@ bool Device::getDeviceIDs(cl_device_type deviceType, uint32_t numEntries, cl_dev
   return true;
 }
 
+bool Device::UpdateStackSize(uint64_t stackSize) {
+  uint32_t maxMemPerThread = info().localMemSizePerCU_ / info().maxThreadsPerCU_;
+  if (maxMemPerThread < stackSize) {
+    return false;
+  }
+  stack_size_ = stackSize;
+  return true;
+}
+
 char* Device::getExtensionString() {
   std::stringstream extStream;
   size_t size;

projects/clr/rocclr/device/device.hpp

@@ -1903,6 +1903,12 @@ class Device : public RuntimeObject {
     return (heap_buffer_ != nullptr) ? heap_buffer_->getDeviceMemory(*this) : nullptr;
   }
 
+  //! Returns stack size set for the device
+  uint64_t StackSize() const { return stack_size_; }
+
+  //! Sets the stack size of the device
+  bool UpdateStackSize(uint64_t stackSize);
+
   //! Does this device allow P2P access?
   bool P2PAccessAllowed() const { return (p2p_access_devices_.size() > 0) ? true : false; }
 
@@ -1962,6 +1968,7 @@ class Device : public RuntimeObject {
   std::once_flag heap_initialized_; //!< Heap buffer initialization flag
   amd::Memory* heap_buffer_;      //!< Preallocated heap buffer for memory allocations on device
   amd::Memory* arena_mem_obj_;    //!< Arena memory object
+  uint64_t stack_size_{0};        //!< Device stack size
 
  private:
   const Isa *isa_;                //!< Device isa

projects/clr/rocclr/device/devkernel.hpp

@@ -498,6 +498,7 @@ class Kernel : public amd::HeapObject {
 
   const uint32_t WorkitemPrivateSegmentByteSize() const { return workitemPrivateSegmentByteSize_; }
   void SetWorkitemPrivateSegmentByteSize(uint32_t size) { workitemPrivateSegmentByteSize_ = size; }
+  const bool KernalHasDynamicCallStack() const { return kernelHasDynamicCallStack_; }
 
   const uint32_t KernargSegmentByteSize() const { return kernargSegmentByteSize_; }
   void SetKernargSegmentByteSize(uint32_t size) { kernargSegmentByteSize_ = size; }
@@ -564,6 +565,7 @@ class Kernel : public amd::HeapObject {
   uint32_t workitemPrivateSegmentByteSize_ = 0;
   uint32_t kernargSegmentByteSize_ = 0;   //!< Size of kernel argument buffer
   uint32_t kernargSegmentAlignment_ = 0;
+  bool kernelHasDynamicCallStack_ = 0;
 
   union Flags {
     struct {

projects/clr/rocclr/device/rocm/rockernel.cpp

@@ -90,6 +90,13 @@ bool LightningKernel::postLoad() {
     return false;
   }
 
+  hsaStatus = hsa_executable_symbol_get_info(symbol, HSA_EXECUTABLE_SYMBOL_INFO_KERNEL_DYNAMIC_CALLSTACK,
+                                             &kernelHasDynamicCallStack_);
+  if (hsaStatus != HSA_STATUS_SUCCESS) {
+    DevLogPrintfError(" Cannot Get Dynamic callstack info, failed with hsa_status: %d \n ", hsaStatus);
+    return false;
+  }
+
   if (!RuntimeHandle().empty()) {
     hsa_executable_symbol_t kernelSymbol;
     int                     variable_size;
@@ -155,7 +162,7 @@ bool LightningKernel::postLoad() {
   workGroupInfo_.localMemSize_ = workgroupGroupSegmentByteSize_;
   workGroupInfo_.usedLDSSize_ = workgroupGroupSegmentByteSize_;
   workGroupInfo_.preferredSizeMultiple_ = wavefront_size;
-  workGroupInfo_.usedStackSize_ = 0;
+  workGroupInfo_.usedStackSize_ = kernelHasDynamicCallStack_;
   workGroupInfo_.wavefrontPerSIMD_ = program()->rocDevice().info().maxWorkItemSizes_[0] / wavefront_size;
   workGroupInfo_.wavefrontSize_ = wavefront_size;
   if (workGroupInfo_.size_ == 0) {

projects/clr/rocclr/device/rocm/rocvirtual.cpp

@@ -2963,6 +2963,14 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
     dispatchPacket.group_segment_size = ldsUsage + sharedMemBytes;
     dispatchPacket.private_segment_size = devKernel->workGroupInfo()->privateMemSize_;
 
+    if ((devKernel->workGroupInfo()->usedStackSize_ & 0x1) == 0x1) {
+      dispatchPacket.private_segment_size += dev().StackSize();
+      uint32_t maxMemPerThread = device().info().localMemSizePerCU_ / device().info().maxThreadsPerCU_;
+      if (dispatchPacket.private_segment_size > maxMemPerThread) {
+        dispatchPacket.private_segment_size = maxMemPerThread;
+      }
+    }
+
     // Pass the header accordingly
     auto aqlHeaderWithOrder = aqlHeader_;
     if (vcmd != nullptr && vcmd->getAnyOrderLaunchFlag()) {
@@ -2970,9 +2978,9 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const
       aqlHeaderWithOrder &= kAqlHeaderMask;
     }
 
-    if (addSystemScope_ || (vcmd != nullptr && 
+    if (addSystemScope_ || (vcmd != nullptr &&
         vcmd->getEventScope() == amd::Device::kCacheStateSystem)) {
-      aqlHeaderWithOrder &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE | 
+      aqlHeaderWithOrder &= ~(HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE |
                               HSA_FENCE_SCOPE_AGENT << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);
       aqlHeaderWithOrder |= (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE |
                              HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE);