diff --git a/projects/clr/rocclr/device/device.cpp b/projects/clr/rocclr/device/device.cpp index d89f526d58..f6417cfaf9 100644 --- a/projects/clr/rocclr/device/device.cpp +++ b/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; diff --git a/projects/clr/rocclr/device/device.hpp b/projects/clr/rocclr/device/device.hpp index 9fab651d24..48510d1eef 100644 --- a/projects/clr/rocclr/device/device.hpp +++ b/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 diff --git a/projects/clr/rocclr/device/devkernel.hpp b/projects/clr/rocclr/device/devkernel.hpp index 29398146db..19b3212dd8 100644 --- a/projects/clr/rocclr/device/devkernel.hpp +++ b/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 { diff --git a/projects/clr/rocclr/device/rocm/rockernel.cpp b/projects/clr/rocclr/device/rocm/rockernel.cpp index f4efc04915..76ab76adcd 100644 --- a/projects/clr/rocclr/device/rocm/rockernel.cpp +++ b/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) { diff --git a/projects/clr/rocclr/device/rocm/rocvirtual.cpp b/projects/clr/rocclr/device/rocm/rocvirtual.cpp index a1cb7c2c1f..3f25f31fd2 100644 --- a/projects/clr/rocclr/device/rocm/rocvirtual.cpp +++ b/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);