SWDEV-531678 - Remove split path from the dispatch (#283)
The split path for blit kernels are no longer necessary, since the new blit kernels don't use the copy size as the global workload
This commit is contained in:
Andryeyev, German
GitHub
@@ -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<const HSAILKernel&>(*(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<size_t>(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<uint32_t>(static_cast<uint32_t>(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<amd::NDRangeContainer&>(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<amd::Memory* const*>(parameters + kernelParams.memoryObjOffset());
|
||||
for (uint32_t i = 0; i < kernel.signature().numMemories(); ++i) {
|
||||
if (nativeMem) {
|
||||
Memory* gpuMem = reinterpret_cast<Memory* const*>(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<uint32_t>(static_cast<uint32_t>(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
|
||||
|
||||
@@ -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<size_t>(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<amd::Memory* const*>(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<address>(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<uint32_t>(sizes.global()[i]);
|
||||
local[i] = static_cast<uint16_t>(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<address>(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<uintptr_t>(printfDbg()->dbgBuffer());
|
||||
if (printfEnabled && bufferPtr) {
|
||||
WriteAqlArgAt(hidden_arguments, bufferPtr, it.size_, it.offset_);
|
||||
}
|
||||
case amd::KernelParameterDescriptor::HiddenPrintfBuffer: {
|
||||
uintptr_t bufferPtr = reinterpret_cast<uintptr_t>(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<uintptr_t>(
|
||||
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<uintptr_t>(
|
||||
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<AmdAqlWrap*>(
|
||||
reinterpret_cast<uint64_t>(schedulerParam_->getHostMem()) + sizeof(SchedulerParam));
|
||||
memset(wrap, 0, sizeof(AmdAqlWrap));
|
||||
wrap->state = AQL_WRAP_DONE;
|
||||
|
||||
spVA = reinterpret_cast<uint64_t>(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<Device::MGSyncInfo*>(
|
||||
dev().MGSync() + Device::kMGInfoSizePerDevice * dev().index() + Device::kMGSyncDataSize);
|
||||
// Update sync data address. Use the offset adjustment to the right location
|
||||
syncInfo->mgs = reinterpret_cast<Device::MGSyncData*>(dev().MGSync() +
|
||||
Device::kMGInfoSizePerDevice * vcmd->firstDevice());
|
||||
} else if (singleGridSync) {
|
||||
syncInfo = reinterpret_cast<Device::MGSyncInfo*>(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<uint64_t>(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<Device&>(dev()).HiddenHeapAlloc(*this);
|
||||
}
|
||||
if (dev().HeapBuffer() != nullptr) {
|
||||
// Initialize hidden heap buffer
|
||||
if (!isGraphCapture) {
|
||||
const_cast<Device&>(dev()).HiddenHeapInit(*this);
|
||||
}
|
||||
// Add heap pointer to the code
|
||||
size_t heap_ptr = static_cast<size_t>(dev().HeapBuffer()->virtualAddress());
|
||||
WriteAqlArgAt(hidden_arguments, heap_ptr, it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenBlockCountX:
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(newGlobalSize[0] / local[0]),
|
||||
it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenBlockCountY:
|
||||
if (sizes.dimensions() >= 2) {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(newGlobalSize[1] / local[1]),
|
||||
it.size_, it.offset_);
|
||||
WriteAqlArgAt(hidden_arguments, buffer, it.size_, it.offset_);
|
||||
} else {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(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<uint32_t>(newGlobalSize[2] / local[2]),
|
||||
it.size_, it.offset_);
|
||||
} else {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint32_t>(1), it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenGroupSizeX:
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[0]), it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenGroupSizeY:
|
||||
if (sizes.dimensions() >= 2) {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[1]), it.size_, it.offset_);
|
||||
} else {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(1), it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenGroupSizeZ:
|
||||
if (sizes.dimensions() >= 3) {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(local[2]), it.size_, it.offset_);
|
||||
} else {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(1), it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenRemainderX:
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(newGlobalSize[0] % local[0]),
|
||||
it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenRemainderY:
|
||||
if (sizes.dimensions() >= 2) {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(newGlobalSize[1] % local[1]),
|
||||
it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenRemainderZ:
|
||||
if (sizes.dimensions() >= 3) {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(newGlobalSize[2] % local[2]),
|
||||
it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenGridDims:
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(sizes.dimensions()),
|
||||
it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenPrivateBase:
|
||||
WriteAqlArgAt(hidden_arguments,
|
||||
reinterpret_cast<amd_queue_t*>(gpu_queue_)->private_segment_aperture_base_hi,
|
||||
it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenSharedBase:
|
||||
WriteAqlArgAt(hidden_arguments,
|
||||
reinterpret_cast<amd_queue_t*>(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<AmdAqlWrap*>(
|
||||
reinterpret_cast<uint64_t>(schedulerParam_->getHostMem()) + sizeof(SchedulerParam));
|
||||
memset(wrap, 0, sizeof(AmdAqlWrap));
|
||||
wrap->state = AQL_WRAP_DONE;
|
||||
spVA = reinterpret_cast<uint64_t>(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<Device::MGSyncInfo*>(
|
||||
dev().MGSync() + Device::kMGInfoSizePerDevice * dev().index() + Device::kMGSyncDataSize);
|
||||
// Update sync data address. Use the offset adjustment to the right location
|
||||
syncInfo->mgs = reinterpret_cast<Device::MGSyncData*>(dev().MGSync() +
|
||||
Device::kMGInfoSizePerDevice * vcmd->firstDevice());
|
||||
} else if (singleGridSync) {
|
||||
syncInfo = reinterpret_cast<Device::MGSyncInfo*>(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<uint64_t>(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<Device&>(dev()).HiddenHeapAlloc(*this);
|
||||
}
|
||||
if (dev().HeapBuffer() != nullptr) {
|
||||
// Initialize hidden heap buffer
|
||||
if (!isGraphCapture) {
|
||||
const_cast<Device&>(dev()).HiddenHeapInit(*this);
|
||||
}
|
||||
// Add heap pointer to the code
|
||||
size_t heap_ptr = static_cast<size_t>(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<uint16_t>(global[0] % local[0]),
|
||||
it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenRemainderY:
|
||||
if (sizes.dimensions() >= 2) {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(global[1] % local[1]),
|
||||
it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenRemainderZ:
|
||||
if (sizes.dimensions() >= 3) {
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(global[2] % local[2]),
|
||||
it.size_, it.offset_);
|
||||
}
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenGridDims:
|
||||
WriteAqlArgAt(hidden_arguments, static_cast<uint16_t>(sizes.dimensions()),
|
||||
it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenPrivateBase:
|
||||
WriteAqlArgAt(hidden_arguments,
|
||||
reinterpret_cast<amd_queue_t*>(gpu_queue_)->private_segment_aperture_base_hi,
|
||||
it.size_, it.offset_);
|
||||
break;
|
||||
case amd::KernelParameterDescriptor::HiddenSharedBase:
|
||||
WriteAqlArgAt(hidden_arguments,
|
||||
reinterpret_cast<amd_queue_t*>(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<address>(
|
||||
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<address>(
|
||||
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<volatile int*>(dev().info().hdpMemFlushCntl);
|
||||
} else if (kernArgImpl == KernelArgImpl::DeviceKernelArgsReadback &&
|
||||
argSize != 0) {
|
||||
_mm_sfence();
|
||||
*(argBuffer + argSize - 1) = *(parameters + argSize - 1);
|
||||
_mm_mfence();
|
||||
auto kSentinel = *reinterpret_cast<volatile unsigned char*>(
|
||||
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<volatile int*>(dev().info().hdpMemFlushCntl);
|
||||
} else if (kernArgImpl == KernelArgImpl::DeviceKernelArgsReadback && argSize != 0) {
|
||||
_mm_sfence();
|
||||
*(argBuffer + argSize - 1) = *(parameters + argSize - 1);
|
||||
_mm_mfence();
|
||||
auto kSentinel = *reinterpret_cast<volatile unsigned char*>(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<uint64_t>(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<uint64_t>(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;
|
||||
}
|
||||
}
|
||||
|
||||
|
||||
مرجع در شماره جدید
Block a user