diff --git a/rocclr/device/pal/palvirtual.cpp b/rocclr/device/pal/palvirtual.cpp index 15b7bf220a..3bd4e60889 100644 --- a/rocclr/device/pal/palvirtual.cpp +++ b/rocclr/device/pal/palvirtual.cpp @@ -2621,25 +2621,17 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::Kernel& kernel, const_address parameters, bool nativeMem, uint32_t sharedMemBytes, bool anyOrder) { - size_t newOffset[3] = {0, 0, 0}; - size_t newGlobalSize[3] = {0, 0, 0}; state_.anyOrder_ = anyOrder; - int dim = -1; - int iteration = 1; - size_t globalStep = 0; - for (uint i = 0; i < sizes.dimensions(); i++) { - newGlobalSize[i] = sizes.global()[i]; - newOffset[i] = sizes.offset()[i]; - } - // Get the HSA kernel object const HSAILKernel& hsaKernel = static_cast(*(kernel.getDeviceKernel(dev()))); // If RGP capturing is enabled, then start SQTT trace if (rgpCaptureEna()) { size_t newLocalSize[3] = {1, 1, 1}; + size_t newGlobalSize[3] = {0, 0, 0}; for (uint i = 0; i < sizes.dimensions(); i++) { + newGlobalSize[i] = sizes.global()[i]; if (sizes.local()[i] != 0) { newLocalSize[i] = sizes.local()[i]; } @@ -2671,13 +2663,8 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, if (PAL_EMBED_KERNEL_MD) { char buf[256]; - sprintf(buf, - "kernel: %s\n" - "private mem size: %x\n" - "group mem size: %x\n", - hsaKernel.name().c_str(), - hsaKernel.spillSegSize(), - hsaKernel.ldsSize()); + sprintf(buf, "kernel: %s\n private mem size: %x\n group mem size: %x\n", + hsaKernel.name().c_str(), hsaKernel.spillSegSize(), hsaKernel.ldsSize()); iCmd()->CmdCommentString(buf); } @@ -2694,128 +2681,77 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, // Add ISA memory object to the resource tracking list AddKernel(kernel); - // Check if it is blit kernel. If it is, then check if split is needed. - if (hsaKernel.isInternalKernel()) { - // Calculate new group size for each submission - for (uint i = 0; i < sizes.dimensions(); i++) { - if (sizes.global()[i] > static_cast(0xffffffff)) { - dim = i; - iteration = sizes.global()[i] / 0xC0000000 + ((sizes.global()[i] % 0xC0000000) ? 1 : 0); - globalStep = (sizes.global()[i] / sizes.local()[i]) / iteration * sizes.local()[dim]; - break; - } + GpuEvent gpuEvent(queues_[MainEngine]->cmdBufId()); + uint32_t id = gpuEvent.id_; + uint64_t vmParentWrap = 0; + uint32_t aql_index = 0; + // Program the kernel arguments for the GPU execution + hsa_kernel_dispatch_packet_t* aqlPkt = hsaKernel.loadArguments( + *this, kernel, sizes, parameters, ldsSize + sharedMemBytes, vmDefQueue, &vmParentWrap, &aql_index); + assert((nullptr != aqlPkt) && "Couldn't load kernel arguments"); + + // Dynamic call stack size is considered to calculate private segment size and scratch regs + // in LightningKernel::postLoad(). As it is not called during hipModuleLaunchKernel unlike + // hipLaunchKernel/hipLaunchKernelGGL, Updated value is passed to dispatch packet. + size_t privateMemSize = hsaKernel.spillSegSize(); + if ((hsaKernel.workGroupInfo()->usedStackSize_ & 0x1) == 0x1) { + privateMemSize = std::max(static_cast(device().StackSize()), + hsaKernel.workGroupInfo()->scratchRegs_ * sizeof(uint32_t)) ; + // Validate privateMemSize is more than max allowed. + size_t maxStackSize = device().MaxStackSize(); + if (privateMemSize > maxStackSize) { + ClPrint(amd::LOG_INFO, amd::LOG_KERN, + "Scratch size (%zu) exceeds max allowed (%zu) for kernel : %s", + privateMemSize, maxStackSize, hsaKernel.name().c_str()); + LogError("Scratch size exceeds max allowed."); + return false; } } - for (int iter = 0; iter < iteration; ++iter) { - GpuEvent gpuEvent(queues_[MainEngine]->cmdBufId()); - uint32_t id = gpuEvent.id_; - // Reset global size for dimension dim if split is needed - if (dim != -1) { - newOffset[dim] = sizes.offset()[dim] + globalStep * iter; - if (((newOffset[dim] + globalStep) < sizes.global()[dim]) && (iter != (iteration - 1))) { - newGlobalSize[dim] = globalStep; - } else { - newGlobalSize[dim] = sizes.global()[dim] - newOffset[dim]; - } - } + // Set up the dispatch information + Pal::DispatchAqlParams dispatchParam = {}; + dispatchParam.pAqlPacket = aqlPkt; + if (privateMemSize > 0) { + const Device::ScratchBuffer* scratch = dev().scratch(hwRing()); + dispatchParam.scratchAddr = scratch->memObj_->vmAddress(); + dispatchParam.scratchSize = scratch->size_; + dispatchParam.scratchOffset = scratch->offset_; + dispatchParam.workitemPrivateSegmentSize = privateMemSize; + } + dispatchParam.pCpuAqlCode = hsaKernel.cpuAqlCode(); + dispatchParam.hsaQueueVa = hsaQueueMem_->vmAddress(); + if (!hsaKernel.prog().isLC() && hsaKernel.workGroupInfo()->wavesPerSimdHint_ != 0) { + constexpr uint32_t kWavesPerSimdLimit = 4; + dispatchParam.wavesPerSh = kWavesPerSimdLimit * + dev().info().cuPerShaderArray_ * dev().info().simdPerCU_; + } else { + dispatchParam.wavesPerSh = 0; + } + dispatchParam.useAtc = dev().settings().svmFineGrainSystem_ ? true : false; + dispatchParam.kernargSegmentSize = hsaKernel.argsBufferSize(); + dispatchParam.aqlPacketIndex = aql_index; + // Run AQL dispatch in HW + eventBegin(MainEngine); + iCmd()->CmdDispatchAql(dispatchParam); - amd::NDRangeContainer tmpSizes(sizes.dimensions(), &newOffset[0], &newGlobalSize[0], - &(const_cast(sizes).local()[0])); + if (id != gpuEvent.id_) { + LogError("Something is wrong. ID mismatch!\n"); + } + eventEnd(MainEngine, gpuEvent); + AqlPacketUpdateTs(aql_index, gpuEvent); - if (iter > 0) { - // Updates the timestamp values, since a CB flush could occur. - // Resource processing was moved from loadArguments() and - // an extra loop is required. - const amd::KernelParameters& kernelParams = kernel.parameters(); - amd::Memory* const* memories = - reinterpret_cast(parameters + kernelParams.memoryObjOffset()); - for (uint32_t i = 0; i < kernel.signature().numMemories(); ++i) { - if (nativeMem) { - Memory* gpuMem = reinterpret_cast(memories)[i]; - if (gpuMem != nullptr) { - gpuMem->setBusy(*this, gpuEvent); - } - } else { - amd::Memory* mem = memories[i]; - if (mem != nullptr) { - dev().getGpuMemory(mem)->setBusy(*this, gpuEvent); - } - } - } - } + // Execute scheduler for device enqueue + if (hsaKernel.dynamicParallelism()) { + PostDeviceEnqueue(kernel, hsaKernel, gpuDefQueue, vmDefQueue, vmParentWrap, &gpuEvent); + } - uint64_t vmParentWrap = 0; - uint32_t aql_index = 0; - // Program the kernel arguments for the GPU execution - hsa_kernel_dispatch_packet_t* aqlPkt = hsaKernel.loadArguments( - *this, kernel, tmpSizes, parameters, ldsSize + sharedMemBytes, vmDefQueue, &vmParentWrap, &aql_index); - if (nullptr == aqlPkt) { - LogError("Couldn't load kernel arguments"); - return false; - } - // Dynamic call stack size is considered to calculate private segment size and scratch regs - // in LightningKernel::postLoad(). As it is not called during hipModuleLaunchKernel unlike - // hipLaunchKernel/hipLaunchKernelGGL, Updated value is passed to dispatch packet. - size_t privateMemSize = hsaKernel.spillSegSize(); - if ((hsaKernel.workGroupInfo()->usedStackSize_ & 0x1) == 0x1) { - privateMemSize = std::max(static_cast(device().StackSize()), - hsaKernel.workGroupInfo()->scratchRegs_ * sizeof(uint32_t)) ; - // Validate privateMemSize is more than max allowed. - size_t maxStackSize = device().MaxStackSize(); - if (privateMemSize > maxStackSize) { - ClPrint(amd::LOG_INFO, amd::LOG_KERN, - "Scratch size (%zu) exceeds max allowed (%zu) for kernel : %s", - privateMemSize, maxStackSize, hsaKernel.name().c_str()); - LogError("Scratch size exceeds max allowed."); - return false; - } - } + // Update the global GPU event + constexpr bool kNeedFLush = false; + setGpuEvent(gpuEvent, kNeedFLush); - // Set up the dispatch information - Pal::DispatchAqlParams dispatchParam = {}; - dispatchParam.pAqlPacket = aqlPkt; - if (privateMemSize > 0) { - const Device::ScratchBuffer* scratch = dev().scratch(hwRing()); - dispatchParam.scratchAddr = scratch->memObj_->vmAddress(); - dispatchParam.scratchSize = scratch->size_; - dispatchParam.scratchOffset = scratch->offset_; - dispatchParam.workitemPrivateSegmentSize = privateMemSize; - } - dispatchParam.pCpuAqlCode = hsaKernel.cpuAqlCode(); - dispatchParam.hsaQueueVa = hsaQueueMem_->vmAddress(); - if (!hsaKernel.prog().isLC() && hsaKernel.workGroupInfo()->wavesPerSimdHint_ != 0) { - constexpr uint32_t kWavesPerSimdLimit = 4; - dispatchParam.wavesPerSh = kWavesPerSimdLimit * - dev().info().cuPerShaderArray_ * dev().info().simdPerCU_; - } else { - dispatchParam.wavesPerSh = 0; - } - dispatchParam.kernargSegmentSize = hsaKernel.argsBufferSize(); - dispatchParam.aqlPacketIndex = aql_index; - // Run AQL dispatch in HW - eventBegin(MainEngine); - iCmd()->CmdDispatchAql(dispatchParam); - - if (id != gpuEvent.id_) { - LogError("Something is wrong. ID mismatch!\n"); - } - eventEnd(MainEngine, gpuEvent); - AqlPacketUpdateTs(aql_index, gpuEvent); - - // Execute scheduler for device enqueue - if (hsaKernel.dynamicParallelism()) { - PostDeviceEnqueue(kernel, hsaKernel, gpuDefQueue, vmDefQueue, vmParentWrap, &gpuEvent); - } - - // Update the global GPU event - constexpr bool kNeedFLush = false; - setGpuEvent(gpuEvent, kNeedFLush); - - if (printfEnabled && !printfDbgHSA().output(*this, printfEnabled, hsaKernel.printfInfo())) { - LogError("Couldn't read printf data from the buffer!\n"); - return false; - } + if (printfEnabled && !printfDbgHSA().output(*this, printfEnabled, hsaKernel.printfInfo())) { + LogError("Couldn't read printf data from the buffer!\n"); + return false; } // Check if image buffer write back is required diff --git a/rocclr/device/rocm/rocvirtual.cpp b/rocclr/device/rocm/rocvirtual.cpp index 8d323c2e77..f4e5799afc 100644 --- a/rocclr/device/rocm/rocvirtual.cpp +++ b/rocclr/device/rocm/rocvirtual.cpp @@ -3337,362 +3337,306 @@ bool VirtualGPU::submitKernelInternal(const amd::NDRangeContainer& sizes, const amd::KernelSignature& signature = kernel.signature(); const amd::KernelParameters& kernelParams = kernel.parameters(); - size_t newOffset[3] = {0, 0, 0}; - size_t newGlobalSize[3] = {0, 0, 0}; - - int dim = -1; - int iteration = 1; - size_t globalStep = 0; - for (uint i = 0; i < sizes.dimensions(); i++) { - newGlobalSize[i] = sizes.global()[i]; - newOffset[i] = sizes.offset()[i]; - } - - if (gpuKernel.isInternalKernel()) { - // Calculate new group size for each submission - for (uint i = 0; i < sizes.dimensions(); i++) { - if (sizes.global()[i] > static_cast(0xffffffff)) { - dim = i; - iteration = sizes.global()[i] / 0xC0000000 + ((sizes.global()[i] % 0xC0000000) ? 1 : 0); - globalStep = (sizes.global()[i] / sizes.local()[i]) / iteration * sizes.local()[dim]; - break; - } - } - } - amd::Memory* const* memories = reinterpret_cast(parameters + kernelParams.memoryObjOffset()); bool isGraphCapture = command_ != nullptr && command_->getPktCapturingState(); - for (int j = 0; j < iteration; j++) { - // Reset global size for dimension dim if split is needed - if (dim != -1) { - newOffset[dim] = sizes.offset()[dim] + globalStep * j; - if (((newOffset[dim] + globalStep) < sizes.global()[dim]) && (j != (iteration - 1))) { - newGlobalSize[dim] = globalStep; - } else { - newGlobalSize[dim] = sizes.global()[dim] - newOffset[dim]; + + ClPrint(amd::LOG_INFO, amd::LOG_KERN, "ShaderName : %s", gpuKernel.getDemangledName().c_str()); + + amd::NDRange local_size(sizes.local()); + address hidden_arguments = const_cast
(parameters); + // Calculate local size if it wasn't provided + devKernel->FindLocalWorkSize(sizes.dimensions(), sizes.global(), local_size); + + uint16_t local[3] = {1, 1, 1}; + uint32_t global[3] = {1, 1, 1}; + for (uint i = 0; i < sizes.dimensions(); i++) { + global[i] = static_cast(sizes.global()[i]); + local[i] = static_cast(local_size[i]); + } + + // Check if runtime has to setup hidden arguments + for (uint32_t i = signature.numParameters(); i < signature.numParametersAll(); ++i) { + const auto& it = signature.at(i); + switch (it.info_.oclObject_) { + case amd::KernelParameterDescriptor::HiddenNone: + break; + case amd::KernelParameterDescriptor::HiddenGlobalOffsetX: { + WriteAqlArgAt(hidden_arguments, sizes.offset()[0], it.size_, it.offset_); + break; } - } - - ClPrint(amd::LOG_INFO, amd::LOG_KERN, "ShaderName : %s", - gpuKernel.getDemangledName().c_str()); - - amd::NDRange local(sizes.local()); - address hidden_arguments = const_cast
(parameters); - // Calculate local size if it wasn't provided - devKernel->FindLocalWorkSize(sizes.dimensions(), sizes.global(), local); - - // Check if runtime has to setup hidden arguments - for (uint32_t i = signature.numParameters(); i < signature.numParametersAll(); ++i) { - const auto& it = signature.at(i); - switch (it.info_.oclObject_) { - case amd::KernelParameterDescriptor::HiddenNone: - break; - case amd::KernelParameterDescriptor::HiddenGlobalOffsetX: { - WriteAqlArgAt(hidden_arguments, newOffset[0], it.size_, it.offset_); - break; + case amd::KernelParameterDescriptor::HiddenGlobalOffsetY: { + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, sizes.offset()[1], it.size_, it.offset_); } - case amd::KernelParameterDescriptor::HiddenGlobalOffsetY: { - if (sizes.dimensions() >= 2) { - WriteAqlArgAt(hidden_arguments, newOffset[1], it.size_, it.offset_); - } - break; + break; + } + case amd::KernelParameterDescriptor::HiddenGlobalOffsetZ: { + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, sizes.offset()[2], it.size_, it.offset_); } - case amd::KernelParameterDescriptor::HiddenGlobalOffsetZ: { - if (sizes.dimensions() >= 3) { - WriteAqlArgAt(hidden_arguments, newOffset[2], it.size_, it.offset_); - } - break; + break; + } + case amd::KernelParameterDescriptor::HiddenPrintfBuffer: { + uintptr_t bufferPtr = reinterpret_cast(printfDbg()->dbgBuffer()); + if (printfEnabled && bufferPtr) { + WriteAqlArgAt(hidden_arguments, bufferPtr, it.size_, it.offset_); } - case amd::KernelParameterDescriptor::HiddenPrintfBuffer: { - uintptr_t bufferPtr = reinterpret_cast(printfDbg()->dbgBuffer()); - if (printfEnabled && bufferPtr) { - WriteAqlArgAt(hidden_arguments, bufferPtr, it.size_, it.offset_); - } - break; - } - case amd::KernelParameterDescriptor::HiddenHostcallBuffer: { - if (amd::IS_HIP) { - if (dev().info().pcie_atomics_) { - uintptr_t buffer = reinterpret_cast( - roc_device_.getOrCreateHostcallBuffer(gpu_queue_, coopGroups, cuMask_)); - if (!buffer) { - LogError("Kernel expects a hostcall buffer, but none found"); - return false; - } - WriteAqlArgAt(hidden_arguments, buffer, it.size_, it.offset_); - } else { - LogError("Pcie atomics not enabled, hostcall not supported"); + break; + } + case amd::KernelParameterDescriptor::HiddenHostcallBuffer: { + if (amd::IS_HIP) { + if (dev().info().pcie_atomics_) { + uintptr_t buffer = reinterpret_cast( + roc_device_.getOrCreateHostcallBuffer(gpu_queue_, coopGroups, cuMask_)); + if (!buffer) { + LogError("Kernel expects a hostcall buffer, but none found"); return false; } - } - break; - } - case amd::KernelParameterDescriptor::HiddenDefaultQueue: { - uint64_t vqVA = 0; - amd::DeviceQueue* defQueue = kernel.program().context().defDeviceQueue(dev()); - if (nullptr != defQueue && devKernel->dynamicParallelism()) { - if (!createVirtualQueue(defQueue->size()) || !createSchedulerParam()) { - return false; - } - vqVA = getVQVirtualAddress(); - } - WriteAqlArgAt(hidden_arguments, vqVA, it.size_, it.offset_); - break; - } - case amd::KernelParameterDescriptor::HiddenCompletionAction: { - uint64_t spVA = 0; - if (nullptr != schedulerParam_ && devKernel->dynamicParallelism()) { - Memory* schedulerMem = dev().getRocMemory(schedulerParam_); - AmdAqlWrap* wrap = reinterpret_cast( - reinterpret_cast(schedulerParam_->getHostMem()) + sizeof(SchedulerParam)); - memset(wrap, 0, sizeof(AmdAqlWrap)); - wrap->state = AQL_WRAP_DONE; - - spVA = reinterpret_cast(schedulerMem->getDeviceMemory()) + sizeof(SchedulerParam); - } - WriteAqlArgAt(hidden_arguments, spVA, it.size_, it.offset_); - break; - } - case amd::KernelParameterDescriptor::HiddenMultiGridSync: { - bool multiGridSync = (vcmd != nullptr) ? vcmd->cooperativeMultiDeviceGroups() : false; - bool singleGridSync = (vcmd != nullptr) ? vcmd->cooperativeGroups() : false; - Device::MGSyncInfo* syncInfo = nullptr; - if (multiGridSync) { - // Find CPU pointer to the right sync info structure. It should be after MGSyncData - syncInfo = reinterpret_cast( - dev().MGSync() + Device::kMGInfoSizePerDevice * dev().index() + Device::kMGSyncDataSize); - // Update sync data address. Use the offset adjustment to the right location - syncInfo->mgs = reinterpret_cast(dev().MGSync() + - Device::kMGInfoSizePerDevice * vcmd->firstDevice()); - } else if (singleGridSync) { - syncInfo = reinterpret_cast(allocKernArg(Device::kSGInfoSize, 64)); - syncInfo->mgs = nullptr; - } - if (multiGridSync || singleGridSync) { - // Update sync data address. - syncInfo->sgs = {0}; - // Fill rest of sync info fields - syncInfo->grid_id = vcmd->gridId(); - syncInfo->num_grids = vcmd->numGrids(); - syncInfo->prev_sum = vcmd->prevGridSum(); - syncInfo->all_sum = vcmd->allGridSum(); - syncInfo->num_wg = vcmd->numWorkgroups(); - } - // Update GPU address for grid sync info. Use the offset adjustment for the right - // location - WriteAqlArgAt(hidden_arguments, reinterpret_cast(syncInfo), it.size_, - it.offset_); - break; - } - case amd::KernelParameterDescriptor::HiddenHeap: - // Allocate hidden heap for HIP applications only - if ((amd::IS_HIP) && (dev().HeapBuffer() == nullptr)) { - const_cast(dev()).HiddenHeapAlloc(*this); - } - if (dev().HeapBuffer() != nullptr) { - // Initialize hidden heap buffer - if (!isGraphCapture) { - const_cast(dev()).HiddenHeapInit(*this); - } - // Add heap pointer to the code - size_t heap_ptr = static_cast(dev().HeapBuffer()->virtualAddress()); - WriteAqlArgAt(hidden_arguments, heap_ptr, it.size_, it.offset_); - } - break; - case amd::KernelParameterDescriptor::HiddenBlockCountX: - WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[0] / local[0]), - it.size_, it.offset_); - break; - case amd::KernelParameterDescriptor::HiddenBlockCountY: - if (sizes.dimensions() >= 2) { - WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[1] / local[1]), - it.size_, it.offset_); + WriteAqlArgAt(hidden_arguments, buffer, it.size_, it.offset_); } else { - WriteAqlArgAt(hidden_arguments, static_cast(1), it.size_, it.offset_); + LogError("Pcie atomics not enabled, hostcall not supported"); + return false; } - break; - case amd::KernelParameterDescriptor::HiddenBlockCountZ: - if (sizes.dimensions() >= 3) { - WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[2] / local[2]), - it.size_, it.offset_); - } else { - WriteAqlArgAt(hidden_arguments, static_cast(1), it.size_, it.offset_); - } - break; - case amd::KernelParameterDescriptor::HiddenGroupSizeX: - WriteAqlArgAt(hidden_arguments, static_cast(local[0]), it.size_, it.offset_); - break; - case amd::KernelParameterDescriptor::HiddenGroupSizeY: - if (sizes.dimensions() >= 2) { - WriteAqlArgAt(hidden_arguments, static_cast(local[1]), it.size_, it.offset_); - } else { - WriteAqlArgAt(hidden_arguments, static_cast(1), it.size_, it.offset_); - } - break; - case amd::KernelParameterDescriptor::HiddenGroupSizeZ: - if (sizes.dimensions() >= 3) { - WriteAqlArgAt(hidden_arguments, static_cast(local[2]), it.size_, it.offset_); - } else { - WriteAqlArgAt(hidden_arguments, static_cast(1), it.size_, it.offset_); - } - break; - case amd::KernelParameterDescriptor::HiddenRemainderX: - WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[0] % local[0]), - it.size_, it.offset_); - break; - case amd::KernelParameterDescriptor::HiddenRemainderY: - if (sizes.dimensions() >= 2) { - WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[1] % local[1]), - it.size_, it.offset_); - } - break; - case amd::KernelParameterDescriptor::HiddenRemainderZ: - if (sizes.dimensions() >= 3) { - WriteAqlArgAt(hidden_arguments, static_cast(newGlobalSize[2] % local[2]), - it.size_, it.offset_); - } - break; - case amd::KernelParameterDescriptor::HiddenGridDims: - WriteAqlArgAt(hidden_arguments, static_cast(sizes.dimensions()), - it.size_, it.offset_); - break; - case amd::KernelParameterDescriptor::HiddenPrivateBase: - WriteAqlArgAt(hidden_arguments, - reinterpret_cast(gpu_queue_)->private_segment_aperture_base_hi, - it.size_, it.offset_); - break; - case amd::KernelParameterDescriptor::HiddenSharedBase: - WriteAqlArgAt(hidden_arguments, - reinterpret_cast(gpu_queue_)->group_segment_aperture_base_hi, - it.size_, it.offset_); - break; - case amd::KernelParameterDescriptor::HiddenQueuePtr: - WriteAqlArgAt(hidden_arguments, gpu_queue_, it.size_, it.offset_); - break; - case amd::KernelParameterDescriptor::HiddenDynamicLdsSize: - WriteAqlArgAt(hidden_arguments, sharedMemBytes, it.size_, it.offset_); - break; + } + break; } + case amd::KernelParameterDescriptor::HiddenDefaultQueue: { + uint64_t vqVA = 0; + amd::DeviceQueue* defQueue = kernel.program().context().defDeviceQueue(dev()); + if (nullptr != defQueue && devKernel->dynamicParallelism()) { + if (!createVirtualQueue(defQueue->size()) || !createSchedulerParam()) { + return false; + } + vqVA = getVQVirtualAddress(); + } + WriteAqlArgAt(hidden_arguments, vqVA, it.size_, it.offset_); + break; + } + case amd::KernelParameterDescriptor::HiddenCompletionAction: { + uint64_t spVA = 0; + if (nullptr != schedulerParam_ && devKernel->dynamicParallelism()) { + Memory* schedulerMem = dev().getRocMemory(schedulerParam_); + AmdAqlWrap* wrap = reinterpret_cast( + reinterpret_cast(schedulerParam_->getHostMem()) + sizeof(SchedulerParam)); + memset(wrap, 0, sizeof(AmdAqlWrap)); + wrap->state = AQL_WRAP_DONE; + spVA = reinterpret_cast(schedulerMem->getDeviceMemory()) + sizeof(SchedulerParam); + } + WriteAqlArgAt(hidden_arguments, spVA, it.size_, it.offset_); + break; + } + case amd::KernelParameterDescriptor::HiddenMultiGridSync: { + bool multiGridSync = (vcmd != nullptr) ? vcmd->cooperativeMultiDeviceGroups() : false; + bool singleGridSync = (vcmd != nullptr) ? vcmd->cooperativeGroups() : false; + Device::MGSyncInfo* syncInfo = nullptr; + if (multiGridSync) { + // Find CPU pointer to the right sync info structure. It should be after MGSyncData + syncInfo = reinterpret_cast( + dev().MGSync() + Device::kMGInfoSizePerDevice * dev().index() + Device::kMGSyncDataSize); + // Update sync data address. Use the offset adjustment to the right location + syncInfo->mgs = reinterpret_cast(dev().MGSync() + + Device::kMGInfoSizePerDevice * vcmd->firstDevice()); + } else if (singleGridSync) { + syncInfo = reinterpret_cast(allocKernArg(Device::kSGInfoSize, 64)); + syncInfo->mgs = nullptr; + } + if (multiGridSync || singleGridSync) { + // Update sync data address. + syncInfo->sgs = {0}; + // Fill rest of sync info fields + syncInfo->grid_id = vcmd->gridId(); + syncInfo->num_grids = vcmd->numGrids(); + syncInfo->prev_sum = vcmd->prevGridSum(); + syncInfo->all_sum = vcmd->allGridSum(); + syncInfo->num_wg = vcmd->numWorkgroups(); + } + // Update GPU address for grid sync info. Use the offset adjustment for the right + // location + WriteAqlArgAt(hidden_arguments, reinterpret_cast(syncInfo), it.size_, + it.offset_); + break; + } + case amd::KernelParameterDescriptor::HiddenHeap: + // Allocate hidden heap for HIP applications only + if ((amd::IS_HIP) && (dev().HeapBuffer() == nullptr)) { + const_cast(dev()).HiddenHeapAlloc(*this); + } + if (dev().HeapBuffer() != nullptr) { + // Initialize hidden heap buffer + if (!isGraphCapture) { + const_cast(dev()).HiddenHeapInit(*this); + } + // Add heap pointer to the code + size_t heap_ptr = static_cast(dev().HeapBuffer()->virtualAddress()); + WriteAqlArgAt(hidden_arguments, heap_ptr, it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenBlockCountX: + WriteAqlArgAt(hidden_arguments, global[0] / local[0], it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenBlockCountY: + WriteAqlArgAt(hidden_arguments, global[1] / local[1], it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenBlockCountZ: + WriteAqlArgAt(hidden_arguments, global[2] / local[2], it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeX: + WriteAqlArgAt(hidden_arguments, local[0], it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeY: + WriteAqlArgAt(hidden_arguments, local[1], it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenGroupSizeZ: + WriteAqlArgAt(hidden_arguments, local[2], it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenRemainderX: + WriteAqlArgAt(hidden_arguments, static_cast(global[0] % local[0]), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenRemainderY: + if (sizes.dimensions() >= 2) { + WriteAqlArgAt(hidden_arguments, static_cast(global[1] % local[1]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenRemainderZ: + if (sizes.dimensions() >= 3) { + WriteAqlArgAt(hidden_arguments, static_cast(global[2] % local[2]), + it.size_, it.offset_); + } + break; + case amd::KernelParameterDescriptor::HiddenGridDims: + WriteAqlArgAt(hidden_arguments, static_cast(sizes.dimensions()), + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenPrivateBase: + WriteAqlArgAt(hidden_arguments, + reinterpret_cast(gpu_queue_)->private_segment_aperture_base_hi, + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenSharedBase: + WriteAqlArgAt(hidden_arguments, + reinterpret_cast(gpu_queue_)->group_segment_aperture_base_hi, + it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenQueuePtr: + WriteAqlArgAt(hidden_arguments, gpu_queue_, it.size_, it.offset_); + break; + case amd::KernelParameterDescriptor::HiddenDynamicLdsSize: + WriteAqlArgAt(hidden_arguments, sharedMemBytes, it.size_, it.offset_); + break; + } + } + address argBuffer = hidden_arguments; + size_t argSize = std::min(gpuKernel.KernargSegmentByteSize(), signature.paramsSize()); + + // Find all parameters for the current kernel + if (!kernel.parameters().deviceKernelArgs() || gpuKernel.isInternalKernel()) { + // Allocate buffer to hold kernel arguments + if (isGraphCapture) { + argBuffer = command_->getKernArgOffset(gpuKernel.KernargSegmentByteSize(), + gpuKernel.KernargSegmentAlignment()); + command_->SetKernelName(gpuKernel.name()); + } else { + ClPrint(amd::LOG_INFO, amd::LOG_KERN, "KernargSegmentByteSize = %lu " + "KernargSegmentAlignment = %lu", gpuKernel.KernargSegmentByteSize(), + gpuKernel.KernargSegmentAlignment()); + argBuffer = reinterpret_cast
( + allocKernArg(gpuKernel.KernargSegmentByteSize(), + gpuKernel.KernargSegmentAlignment())); } - address argBuffer = hidden_arguments; - size_t argSize = std::min(gpuKernel.KernargSegmentByteSize(), signature.paramsSize()); + nontemporalMemcpy(argBuffer, parameters, argSize); - // Find all parameters for the current kernel - if (!kernel.parameters().deviceKernelArgs() || gpuKernel.isInternalKernel()) { - // Allocate buffer to hold kernel arguments - if (isGraphCapture) { - argBuffer = command_->getKernArgOffset(gpuKernel.KernargSegmentByteSize(), - gpuKernel.KernargSegmentAlignment()); - command_->SetKernelName(gpuKernel.name()); - } else { - ClPrint(amd::LOG_INFO, amd::LOG_KERN, "KernargSegmentByteSize = %lu " - "KernargSegmentAlignment = %lu", gpuKernel.KernargSegmentByteSize(), - gpuKernel.KernargSegmentAlignment()); - argBuffer = reinterpret_cast
( - allocKernArg(gpuKernel.KernargSegmentByteSize(), - gpuKernel.KernargSegmentAlignment())); - } - - nontemporalMemcpy(argBuffer, parameters, argSize); - - if (roc_device_.info().largeBar_ && !isGraphCapture) { - const auto kernArgImpl = dev().settings().kernel_arg_impl_; - - if (kernArgImpl == KernelArgImpl::DeviceKernelArgsHDP) { - *dev().info().hdpMemFlushCntl = 1u; - auto kSentinel = *reinterpret_cast(dev().info().hdpMemFlushCntl); - } else if (kernArgImpl == KernelArgImpl::DeviceKernelArgsReadback && - argSize != 0) { - _mm_sfence(); - *(argBuffer + argSize - 1) = *(parameters + argSize - 1); - _mm_mfence(); - auto kSentinel = *reinterpret_cast( - argBuffer + argSize - 1); - } + if (roc_device_.info().largeBar_ && !isGraphCapture) { + const auto kernArgImpl = dev().settings().kernel_arg_impl_; + if (kernArgImpl == KernelArgImpl::DeviceKernelArgsHDP) { + *dev().info().hdpMemFlushCntl = 1u; + auto kSentinel = *reinterpret_cast(dev().info().hdpMemFlushCntl); + } else if (kernArgImpl == KernelArgImpl::DeviceKernelArgsReadback && argSize != 0) { + _mm_sfence(); + *(argBuffer + argSize - 1) = *(parameters + argSize - 1); + _mm_mfence(); + auto kSentinel = *reinterpret_cast(argBuffer + argSize - 1); } } + } - // Check for group memory overflow - //! @todo Check should be in HSA - here we should have at most an assert - assert(dev().info().localMemSizePerCU_ > 0); - if (ldsUsage > dev().info().localMemSizePerCU_) { - LogError("No local memory available\n"); + // Check for group memory overflow + //! @todo Check should be in HSA - here we should have at most an assert + assert(dev().info().localMemSizePerCU_ > 0); + if (ldsUsage > dev().info().localMemSizePerCU_) { + LogError("No local memory available\n"); + return false; + } + + // Initialize the dispatch Packet + hsa_kernel_dispatch_packet_t dispatchPacket{}; + + dispatchPacket.header = kInvalidAql; + dispatchPacket.kernel_object = gpuKernel.KernelCodeHandle(); + + dispatchPacket.grid_size_x = global[0]; + dispatchPacket.grid_size_y = global[1]; + dispatchPacket.grid_size_z = global[2]; + + dispatchPacket.workgroup_size_x = local[0]; + dispatchPacket.workgroup_size_y = local[1]; + dispatchPacket.workgroup_size_z = local[2]; + + dispatchPacket.kernarg_address = argBuffer; + dispatchPacket.group_segment_size = ldsUsage + sharedMemBytes; + dispatchPacket.private_segment_size = devKernel->workGroupInfo()->privateMemSize_; + if ((devKernel->workGroupInfo()->usedStackSize_ & 0x1) == 0x1) { + dispatchPacket.private_segment_size = std::max(dev().StackSize(), + dispatchPacket.private_segment_size); + // Validate privateMemSize is more than max allowed. + size_t maxStackSize = dev().MaxStackSize(); + if (dispatchPacket.private_segment_size > maxStackSize) { + ClPrint(amd::LOG_INFO, amd::LOG_KERN, + "Scratch size (%u) exceeds max allowed (%zu) for kernel : %s", + dispatchPacket.private_segment_size, maxStackSize, gpuKernel.getDemangledName().c_str()); + LogError("Scratch size exceeds max allowed."); return false; } + } - // Initialize the dispatch Packet - hsa_kernel_dispatch_packet_t dispatchPacket; - memset(&dispatchPacket, 0, sizeof(dispatchPacket)); - - dispatchPacket.header = kInvalidAql; - dispatchPacket.kernel_object = gpuKernel.KernelCodeHandle(); - - // dispatchPacket.header = aqlHeader_; - // dispatchPacket.setup |= sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; - dispatchPacket.grid_size_x = sizes.dimensions() > 0 ? newGlobalSize[0] : 1; - dispatchPacket.grid_size_y = sizes.dimensions() > 1 ? newGlobalSize[1] : 1; - dispatchPacket.grid_size_z = sizes.dimensions() > 2 ? newGlobalSize[2] : 1; - - dispatchPacket.workgroup_size_x = sizes.dimensions() > 0 ? local[0] : 1; - dispatchPacket.workgroup_size_y = sizes.dimensions() > 1 ? local[1] : 1; - dispatchPacket.workgroup_size_z = sizes.dimensions() > 2 ? local[2] : 1; - - dispatchPacket.kernarg_address = argBuffer; - dispatchPacket.group_segment_size = ldsUsage + sharedMemBytes; - dispatchPacket.private_segment_size = devKernel->workGroupInfo()->privateMemSize_; - - if ((devKernel->workGroupInfo()->usedStackSize_ & 0x1) == 0x1) { - dispatchPacket.private_segment_size = std::max(dev().StackSize(), - dispatchPacket.private_segment_size); - // Validate privateMemSize is more than max allowed. - size_t maxStackSize = dev().MaxStackSize(); - if (dispatchPacket.private_segment_size > maxStackSize) { - ClPrint(amd::LOG_INFO, amd::LOG_KERN, - "Scratch size (%u) exceeds max allowed (%zu) for kernel : %s", - dispatchPacket.private_segment_size, maxStackSize, - gpuKernel.getDemangledName().c_str()); - LogError("Scratch size exceeds max allowed."); - return false; - } - } - - // Pass the header accordingly - auto aqlHeaderWithOrder = aqlHeader_; - if (vcmd != nullptr && vcmd->getAnyOrderLaunchFlag()) { + // Pass the header accordingly + auto aqlHeaderWithOrder = aqlHeader_; + if (vcmd != nullptr) { + if (vcmd->getAnyOrderLaunchFlag()) { constexpr uint32_t kAqlHeaderMask = ~(1 << HSA_PACKET_HEADER_BARRIER); aqlHeaderWithOrder &= kAqlHeaderMask; } - - if (vcmd != nullptr && vcmd->getCommandEntryScope() == amd::Device::kCacheStateSystem) { + if (vcmd->getCommandEntryScope() == amd::Device::kCacheStateSystem) { addSystemScope_ = true; } + } - // Copy scheduler's AQL packet for possible relaunch from the scheduler itself - if (aql_packet != nullptr) { - *aql_packet = dispatchPacket; - aql_packet->header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | - (1 << HSA_PACKET_HEADER_BARRIER) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | - (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); - aql_packet->setup = sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; + // Copy scheduler's AQL packet for possible relaunch from the scheduler itself + if (aql_packet != nullptr) { + *aql_packet = dispatchPacket; + aql_packet->header = (HSA_PACKET_TYPE_KERNEL_DISPATCH << HSA_PACKET_HEADER_TYPE) | + (1 << HSA_PACKET_HEADER_BARRIER) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_ACQUIRE_FENCE_SCOPE) | + (HSA_FENCE_SCOPE_SYSTEM << HSA_PACKET_HEADER_RELEASE_FENCE_SCOPE); + aql_packet->setup = sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS; + } + + if (isGraphCapture) { + // Dispatch the packet + if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, + (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), + GPU_FLUSH_ON_EXECUTION, command_->getPktCapturingState(), + command_->getAqlPacket())) { + return false; } - - if (isGraphCapture) { - // Dispatch the packet - if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, - (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), - GPU_FLUSH_ON_EXECUTION, command_->getPktCapturingState(), - command_->getAqlPacket())) { - return false; - } - } else { - if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, - (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), - GPU_FLUSH_ON_EXECUTION, false, nullptr, attach_signal)) { - return false; - } + } else { + if (!dispatchAqlPacket(&dispatchPacket, aqlHeaderWithOrder, + (sizes.dimensions() << HSA_KERNEL_DISPATCH_PACKET_SETUP_DIMENSIONS), + GPU_FLUSH_ON_EXECUTION, false, nullptr, attach_signal)) { + return false; } }