intel/compute-runtime
1
0
Fork 0
mirror of https://github.com/intel/compute-runtime.git synced 2026-01-13 01:42:56 +08:00
Code Issues Packages Projects Releases Wiki Activity

refactor: pass ze group count as a reference, not a pointer

Browse Source 
Signed-off-by: Mateusz Jablonski <mateusz.jablonski@intel.com>
This commit is contained in:
Mateusz Jablonski
2023-09-27 16:27:47 +00:00
committed by Compute-Runtime-Automation
parent 2aaf5a1f03
commit cb730d11f4
39 changed files with 551 additions and 568 deletions
Download Patch File Download Diff File Expand all files Collapse all files

8
level_zero/api/core/ze_module_api_entrypoints.h
View File

@@ -152,7 +152,7 @@ ze_result_t zeCommandListAppendLaunchKernel(
L0::CmdListKernelLaunchParams launchParams = {};
launchParams.skipInOrderNonWalkerSignaling = cmdList->skipInOrderNonWalkerSignalingAllowed(hSignalEvent);
return cmdList->appendLaunchKernel(kernelHandle, launchKernelArgs, hSignalEvent, numWaitEvents, phWaitEvents, launchParams, false);
return cmdList->appendLaunchKernel(kernelHandle, *launchKernelArgs, hSignalEvent, numWaitEvents, phWaitEvents, launchParams, false);
}
ze_result_t zeCommandListAppendLaunchCooperativeKernel(
@@ -162,7 +162,7 @@ ze_result_t zeCommandListAppendLaunchCooperativeKernel(
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
return L0::CommandList::fromHandle(hCommandList)->appendLaunchCooperativeKernel(kernelHandle, launchKernelArgs, hSignalEvent, numWaitEvents, phWaitEvents, false);
return L0::CommandList::fromHandle(hCommandList)->appendLaunchCooperativeKernel(kernelHandle, *launchKernelArgs, hSignalEvent, numWaitEvents, phWaitEvents, false);
}
ze_result_t zeCommandListAppendLaunchKernelIndirect(
@@ -172,7 +172,7 @@ ze_result_t zeCommandListAppendLaunchKernelIndirect(
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
return L0::CommandList::fromHandle(hCommandList)->appendLaunchKernelIndirect(kernelHandle, pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents, false);
return L0::CommandList::fromHandle(hCommandList)->appendLaunchKernelIndirect(kernelHandle, *pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents, false);
}
ze_result_t zeCommandListAppendLaunchMultipleKernelsIndirect(
@@ -184,7 +184,7 @@ ze_result_t zeCommandListAppendLaunchMultipleKernelsIndirect(
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents) {
return L0::CommandList::fromHandle(hCommandList)->appendLaunchMultipleKernelsIndirect(numKernels, kernelHandles, pCountBuffer, pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents, false);
return L0::CommandList::fromHandle(hCommandList)->appendLaunchMultipleKernelsIndirect(numKernels, kernelHandles, pCountBuffer, *pLaunchArgumentsBuffer, hSignalEvent, numWaitEvents, phWaitEvents, false);
}
ze_result_t zeKernelGetName(

8
level_zero/core/source/cmdlist/cmdlist.h
View File

@@ -104,21 +104,21 @@ struct CommandList : _ze_command_list_handle_t {
virtual ze_result_t appendImageCopy(ze_image_handle_t hDstImage, ze_image_handle_t hSrcImage,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) = 0;
virtual ze_result_t appendLaunchKernel(ze_kernel_handle_t kernelHandle, const ze_group_count_t *threadGroupDimensions,
virtual ze_result_t appendLaunchKernel(ze_kernel_handle_t kernelHandle, const ze_group_count_t &threadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents,
const CmdListKernelLaunchParams &launchParams, bool relaxedOrderingDispatch) = 0;
virtual ze_result_t appendLaunchCooperativeKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *launchKernelArgs,
const ze_group_count_t &launchKernelArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *waitEventHandles, bool relaxedOrderingDispatch) = 0;
virtual ze_result_t appendLaunchKernelIndirect(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *pDispatchArgumentsBuffer,
const ze_group_count_t &pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) = 0;
virtual ze_result_t appendLaunchMultipleKernelsIndirect(uint32_t numKernels, const ze_kernel_handle_t *kernelHandles,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer, ze_event_handle_t hEvent,
const ze_group_count_t &pLaunchArgumentsBuffer, ze_event_handle_t hEvent,
uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) = 0;
virtual ze_result_t appendMemAdvise(ze_device_handle_t hDevice, const void *ptr, size_t size,
ze_memory_advice_t advice) = 0;

20
level_zero/core/source/cmdlist/cmdlist_hw.h
View File

@@ -102,23 +102,23 @@ struct CommandListCoreFamily : CommandListImp {
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) override;
ze_result_t appendLaunchKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents,
const CmdListKernelLaunchParams &launchParams, bool relaxedOrderingDispatch) override;
ze_result_t appendLaunchCooperativeKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *launchKernelArgs,
const ze_group_count_t &launchKernelArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *waitEventHandles, bool relaxedOrderingDispatch) override;
ze_result_t appendLaunchKernelIndirect(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *pDispatchArgumentsBuffer,
const ze_group_count_t &pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) override;
ze_result_t appendLaunchMultipleKernelsIndirect(uint32_t numKernels,
const ze_kernel_handle_t *kernelHandles,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
const ze_group_count_t &pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) override;
@@ -248,11 +248,11 @@ struct CommandListCoreFamily : CommandListImp {
Event *signalEvent);
MOCKABLE_VIRTUAL ze_result_t appendLaunchKernelWithParams(Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
Event *event,
const CmdListKernelLaunchParams &launchParams);
MOCKABLE_VIRTUAL ze_result_t appendLaunchKernelSplit(Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
Event *event,
const CmdListKernelLaunchParams &launchParams);
@@ -266,9 +266,9 @@ struct CommandListCoreFamily : CommandListImp {
void appendWaitOnSingleEvent(Event *event, bool relaxedOrderingAllowed);
ze_result_t prepareIndirectParams(const ze_group_count_t *threadGroupDimensions);
void updateStreamPropertiesForRegularCommandLists(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect);
void updateStreamPropertiesForFlushTaskDispatchFlags(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect);
void updateStreamProperties(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect);
void updateStreamPropertiesForRegularCommandLists(Kernel &kernel, bool isCooperative, const ze_group_count_t &threadGroupDimensions, bool isIndirect);
void updateStreamPropertiesForFlushTaskDispatchFlags(Kernel &kernel, bool isCooperative, const ze_group_count_t &threadGroupDimensions, bool isIndirect);
void updateStreamProperties(Kernel &kernel, bool isCooperative, const ze_group_count_t &threadGroupDimensions, bool isIndirect);
void clearCommandsToPatch();
size_t getTotalSizeForCopyRegion(const ze_copy_region_t *region, uint32_t pitch, uint32_t slicePitch);
@@ -279,7 +279,7 @@ struct CommandListCoreFamily : CommandListImp {
const void **pRanges);
ze_result_t setGlobalWorkSizeIndirect(NEO::CrossThreadDataOffset offsets[3], uint64_t crossThreadAddress, uint32_t lws[3]);
ze_result_t programSyncBuffer(Kernel &kernel, NEO::Device &device, const ze_group_count_t *threadGroupDimensions);
ze_result_t programSyncBuffer(Kernel &kernel, NEO::Device &device, const ze_group_count_t &threadGroupDimensions);
void appendWriteKernelTimestamp(Event *event, bool beforeWalker, bool maskLsb, bool workloadPartition);
void adjustWriteKernelTimestamp(uint64_t globalAddress, uint64_t contextAddress, bool maskLsb, uint32_t mask, bool workloadPartition);
void appendEventForProfiling(Event *event, bool beforeWalker, bool skipBarrierForEndProfiling);

67
level_zero/core/source/cmdlist/cmdlist_hw.inl
View File

@@ -321,7 +321,7 @@ void CommandListCoreFamily<gfxCoreFamily>::programL3(bool isSLMused) {}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents,
@@ -372,7 +372,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(ze_kernel_h
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchCooperativeKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *launchKernelArgs,
const ze_group_count_t &launchKernelArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *waitEventHandles, bool relaxedOrderingDispatch) {
@@ -403,7 +403,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchCooperativeKernel(
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelIndirect(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *pDispatchArgumentsBuffer,
const ze_group_count_t &pDispatchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) {
@@ -441,7 +441,7 @@ template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchMultipleKernelsIndirect(uint32_t numKernels,
const ze_kernel_handle_t *kernelHandles,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
const ze_group_count_t &pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) {
@@ -462,7 +462,6 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchMultipleKernelsInd
}
appendEventForProfiling(event, true, false);
const bool haveLaunchArguments = pLaunchArgumentsBuffer != nullptr;
auto allocData = device->getDriverHandle()->getSvmAllocsManager()->getSVMAlloc(pNumLaunchArguments);
auto alloc = allocData->gpuAllocations.getGraphicsAllocation(device->getRootDeviceIndex());
commandContainer.addToResidencyContainer(alloc);
@@ -471,7 +470,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchMultipleKernelsInd
NEO::EncodeMathMMIO<GfxFamily>::encodeGreaterThanPredicate(commandContainer, alloc->getGpuAddress(), i);
ret = appendLaunchKernelWithParams(Kernel::fromHandle(kernelHandles[i]),
haveLaunchArguments ? &pLaunchArgumentsBuffer[i] : nullptr,
pLaunchArgumentsBuffer,
nullptr, launchParams);
if (ret) {
return ret;
@@ -709,7 +708,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyFromMemory(ze_i
CmdListKernelLaunchParams launchParams = {};
launchParams.isBuiltInKernel = true;
auto status = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(), &kernelArgs,
auto status = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(), kernelArgs,
event, numWaitEvents, phWaitEvents,
launchParams, relaxedOrderingDispatch);
addToMappedEventList(Event::fromHandle(hEvent));
@@ -867,7 +866,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyToMemory(void *
launchParams.isDestinationAllocationInSystemMemory =
(dstAllocationType == NEO::AllocationType::BUFFER_HOST_MEMORY) ||
(dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR);
ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(), &kernelArgs,
ret = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(), kernelArgs,
event, numWaitEvents, phWaitEvents, launchParams, relaxedOrderingDispatch);
addToMappedEventList(event);
@@ -1012,7 +1011,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendImageCopyRegion(ze_image
CmdListKernelLaunchParams launchParams = {};
launchParams.isBuiltInKernel = true;
auto status = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(kernel->toHandle(), &kernelArgs,
auto status = CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(kernel->toHandle(), kernelArgs,
event, numWaitEvents, phWaitEvents,
launchParams, relaxedOrderingDispatch);
addToMappedEventList(event);
@@ -1149,7 +1148,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryCopyKernelWithGA(v
(dstAllocationType == NEO::AllocationType::SVM_CPU) ||
(dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelSplit(builtinKernel, &dispatchKernelArgs, signalEvent, launchParams);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelSplit(builtinKernel, dispatchKernelArgs, signalEvent, launchParams);
}
template <GFXCORE_FAMILY gfxCoreFamily>
@@ -1647,7 +1646,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryCopyKernel3d(Align
launchParams.isDestinationAllocationInSystemMemory =
(dstAllocationType == NEO::AllocationType::BUFFER_HOST_MEMORY) ||
(dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(), &dispatchKernelArgs, signalEvent, numWaitEvents,
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(), dispatchKernelArgs, signalEvent, numWaitEvents,
phWaitEvents, launchParams, relaxedOrderingDispatch);
}
@@ -1715,7 +1714,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryCopyKernel2d(Align
(dstAllocationType == NEO::AllocationType::BUFFER_HOST_MEMORY) ||
(dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR);
return CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernel(builtinKernel->toHandle(),
&dispatchKernelArgs, signalEvent,
dispatchKernelArgs, signalEvent,
numWaitEvents,
phWaitEvents,
launchParams, relaxedOrderingDispatch);
@@ -1749,7 +1748,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendUnalignedFillKernel(bool
builtinKernel->setArgumentValue(1, sizeof(dstAllocation.offset), &dstAllocation.offset);
builtinKernel->setArgumentValue(2, sizeof(value), &value);
auto res = appendLaunchKernelSplit(builtinKernel, &dispatchKernelRemainderArgs, signalEvent, launchParams);
auto res = appendLaunchKernelSplit(builtinKernel, dispatchKernelRemainderArgs, signalEvent, launchParams);
if (res) {
return res;
}
@@ -1872,7 +1871,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryFill(void *ptr,
builtinKernel->setArgumentValue(1, sizeof(fillArguments.mainOffset), &fillArguments.mainOffset);
builtinKernel->setArgumentValue(2, sizeof(value), &value);
res = appendLaunchKernelSplit(builtinKernel, &dispatchKernelArgs, signalEvent, launchParams);
res = appendLaunchKernelSplit(builtinKernel, dispatchKernelArgs, signalEvent, launchParams);
if (res) {
return res;
}
@@ -1917,7 +1916,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryFill(void *ptr,
builtinKernel->setArgumentValue(3, sizeof(fillArguments.patternSizeInEls), &fillArguments.patternSizeInEls);
ze_group_count_t dispatchKernelArgs{static_cast<uint32_t>(fillArguments.groups), 1u, 1u};
res = appendLaunchKernelSplit(builtinKernel, &dispatchKernelArgs, signalEvent, launchParams);
res = appendLaunchKernelSplit(builtinKernel, dispatchKernelArgs, signalEvent, launchParams);
if (res) {
return res;
}
@@ -1945,7 +1944,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryFill(void *ptr,
patternGfxAlloc, nullptr);
builtinKernelRemainder->setArgumentValue(3, sizeof(patternAllocationSize), &patternAllocationSize);
res = appendLaunchKernelSplit(builtinKernelRemainder, &dispatchKernelArgs, signalEvent, launchParams);
res = appendLaunchKernelSplit(builtinKernelRemainder, dispatchKernelArgs, signalEvent, launchParams);
if (res) {
return res;
}
@@ -1976,7 +1975,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendMemoryFill(void *ptr,
patternGfxAlloc, nullptr);
builtinKernelRemainder->setArgumentValue(3, sizeof(patternAllocationSize), &patternAllocationSize);
res = appendLaunchKernelSplit(builtinKernelRemainder, &dispatchKernelArgs, signalEvent, launchParams);
res = appendLaunchKernelSplit(builtinKernelRemainder, dispatchKernelArgs, signalEvent, launchParams);
if (res) {
return res;
}
@@ -2439,13 +2438,13 @@ void CommandListCoreFamily<gfxCoreFamily>::appendSignalInOrderDependencyCounter(
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::programSyncBuffer(Kernel &kernel, NEO::Device &device,
const ze_group_count_t *threadGroupDimensions) {
const ze_group_count_t &threadGroupDimensions) {
uint32_t maximalNumberOfWorkgroupsAllowed;
auto ret = kernel.suggestMaxCooperativeGroupCount(&maximalNumberOfWorkgroupsAllowed, this->engineGroupType,
device.isEngineInstanced());
UNRECOVERABLE_IF(ret != ZE_RESULT_SUCCESS);
size_t requestedNumberOfWorkgroups = (threadGroupDimensions->groupCountX * threadGroupDimensions->groupCountY *
threadGroupDimensions->groupCountZ);
size_t requestedNumberOfWorkgroups = (threadGroupDimensions.groupCountX * threadGroupDimensions.groupCountY *
threadGroupDimensions.groupCountZ);
if (requestedNumberOfWorkgroups > maximalNumberOfWorkgroupsAllowed) {
return ZE_RESULT_ERROR_INVALID_ARGUMENT;
}
@@ -2681,7 +2680,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendQueryKernelTimestamps(
launchParams.isDestinationAllocationInSystemMemory =
(dstAllocationType == NEO::AllocationType::BUFFER_HOST_MEMORY) ||
(dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR);
auto appendResult = appendLaunchKernel(builtinKernel->toHandle(), &dispatchKernelArgs, hSignalEvent, numWaitEvents,
auto appendResult = appendLaunchKernel(builtinKernel->toHandle(), dispatchKernelArgs, hSignalEvent, numWaitEvents,
phWaitEvents, launchParams, false);
if (appendResult != ZE_RESULT_SUCCESS) {
return appendResult;
@@ -2737,7 +2736,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::prepareIndirectParams(const ze
}
template <GFXCORE_FAMILY gfxCoreFamily>
void CommandListCoreFamily<gfxCoreFamily>::updateStreamProperties(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect) {
void CommandListCoreFamily<gfxCoreFamily>::updateStreamProperties(Kernel &kernel, bool isCooperative, const ze_group_count_t &threadGroupDimensions, bool isIndirect) {
if (this->isFlushTaskSubmissionEnabled) {
updateStreamPropertiesForFlushTaskDispatchFlags(kernel, isCooperative, threadGroupDimensions, isIndirect);
} else {
@@ -2746,29 +2745,27 @@ void CommandListCoreFamily<gfxCoreFamily>::updateStreamProperties(Kernel &kernel
}
template <GFXCORE_FAMILY gfxCoreFamily>
inline bool getFusedEuDisabled(Kernel &kernel, Device *device, const ze_group_count_t *threadGroupDimensions, bool isIndirect) {
inline bool getFusedEuDisabled(Kernel &kernel, Device *device, const ze_group_count_t &threadGroupDimensions, bool isIndirect) {
auto &kernelAttributes = kernel.getKernelDescriptor().kernelAttributes;
bool fusedEuDisabled = kernelAttributes.flags.requiresDisabledEUFusion;
if (static_cast<DeviceImp *>(device)->calculationForDisablingEuFusionWithDpasNeeded) {
auto &productHelper = device->getProductHelper();
if (threadGroupDimensions) {
uint32_t *groupCountPtr = nullptr;
uint32_t groupCount[3] = {};
if (!isIndirect) {
groupCount[0] = threadGroupDimensions->groupCountX;
groupCount[1] = threadGroupDimensions->groupCountY;
groupCount[2] = threadGroupDimensions->groupCountZ;
groupCountPtr = groupCount;
}
fusedEuDisabled |= productHelper.isFusedEuDisabledForDpas(kernelAttributes.flags.usesSystolicPipelineSelectMode, kernel.getGroupSize(), groupCountPtr, device->getHwInfo());
uint32_t *groupCountPtr = nullptr;
uint32_t groupCount[3] = {};
if (!isIndirect) {
groupCount[0] = threadGroupDimensions.groupCountX;
groupCount[1] = threadGroupDimensions.groupCountY;
groupCount[2] = threadGroupDimensions.groupCountZ;
groupCountPtr = groupCount;
}
fusedEuDisabled |= productHelper.isFusedEuDisabledForDpas(kernelAttributes.flags.usesSystolicPipelineSelectMode, kernel.getGroupSize(), groupCountPtr, device->getHwInfo());
}
return fusedEuDisabled;
}
template <GFXCORE_FAMILY gfxCoreFamily>
void CommandListCoreFamily<gfxCoreFamily>::updateStreamPropertiesForFlushTaskDispatchFlags(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect) {
void CommandListCoreFamily<gfxCoreFamily>::updateStreamPropertiesForFlushTaskDispatchFlags(Kernel &kernel, bool isCooperative, const ze_group_count_t &threadGroupDimensions, bool isIndirect) {
auto &kernelAttributes = kernel.getKernelDescriptor().kernelAttributes;
bool fusedEuDisabled = getFusedEuDisabled<gfxCoreFamily>(kernel, this->device, threadGroupDimensions, isIndirect);
@@ -2785,7 +2782,7 @@ void CommandListCoreFamily<gfxCoreFamily>::updateStreamPropertiesForFlushTaskDis
}
template <GFXCORE_FAMILY gfxCoreFamily>
void CommandListCoreFamily<gfxCoreFamily>::updateStreamPropertiesForRegularCommandLists(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect) {
void CommandListCoreFamily<gfxCoreFamily>::updateStreamPropertiesForRegularCommandLists(Kernel &kernel, bool isCooperative, const ze_group_count_t &threadGroupDimensions, bool isIndirect) {
using VFE_STATE_TYPE = typename GfxFamily::VFE_STATE_TYPE;
size_t currentSurfaceStateSize = NEO::StreamPropertySizeT::initValue;

6
level_zero/core/source/cmdlist/cmdlist_hw_immediate.h
View File

@@ -52,13 +52,13 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily<gfxCoreFami
CommandListCoreFamilyImmediate(uint32_t numIddsPerBlock);
ze_result_t appendLaunchKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents,
const CmdListKernelLaunchParams &launchParams, bool relaxedOrderingDispatch) override;
ze_result_t appendLaunchKernelIndirect(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *pDispatchArgumentsBuffer,
const ze_group_count_t &pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) override;
@@ -144,7 +144,7 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily<gfxCoreFami
ze_event_handle_t *phWaitEvents) override;
ze_result_t appendLaunchCooperativeKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *launchKernelArgs,
const ze_group_count_t &launchKernelArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *waitEventHandles, bool relaxedOrderingDispatch) override;

6
level_zero/core/source/cmdlist/cmdlist_hw_immediate.inl
View File

@@ -414,7 +414,7 @@ bool CommandListCoreFamilyImmediate<gfxCoreFamily>::skipInOrderNonWalkerSignalin
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchKernel(
ze_kernel_handle_t kernelHandle, const ze_group_count_t *threadGroupDimensions,
ze_kernel_handle_t kernelHandle, const ze_group_count_t &threadGroupDimensions,
ze_event_handle_t hSignalEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents,
const CmdListKernelLaunchParams &launchParams, bool relaxedOrderingDispatch) {
@@ -465,7 +465,7 @@ void CommandListCoreFamilyImmediate<gfxCoreFamily>::handleInOrderNonWalkerSignal
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchKernelIndirect(
ze_kernel_handle_t kernelHandle, const ze_group_count_t *pDispatchArgumentsBuffer,
ze_kernel_handle_t kernelHandle, const ze_group_count_t &pDispatchArgumentsBuffer,
ze_event_handle_t hSignalEvent, uint32_t numWaitEvents, ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) {
relaxedOrderingDispatch = isRelaxedOrderingDispatchAllowed(numWaitEvents);
@@ -848,7 +848,7 @@ ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendMemoryRangesBar
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamilyImmediate<gfxCoreFamily>::appendLaunchCooperativeKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *launchKernelArgs,
const ze_group_count_t &launchKernelArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *waitEventHandles, bool relaxedOrderingDispatch) {

18
level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl
View File

@@ -47,7 +47,7 @@ bool CommandListCoreFamily<gfxCoreFamily>::isInOrderNonWalkerSignalingRequired(c
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
Event *event,
const CmdListKernelLaunchParams &launchParams) {
UNRECOVERABLE_IF(kernel == nullptr);
@@ -66,7 +66,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
DBG_LOG(PrintDispatchParameters, "Kernel: ", kernelInfo->kernelDescriptor.kernelMetadata.kernelName,
", Group size: ", kernel->getGroupSize()[0], ", ", kernel->getGroupSize()[1], ", ", kernel->getGroupSize()[2],
", Group count: ", threadGroupDimensions->groupCountX, ", ", threadGroupDimensions->groupCountY, ", ", threadGroupDimensions->groupCountZ,
", Group count: ", threadGroupDimensions.groupCountX, ", ", threadGroupDimensions.groupCountY, ", ", threadGroupDimensions.groupCountZ,
", SIMD: ", kernelInfo->getMaxSimdSize());
if (this->immediateCmdListHeapSharing || this->stateBaseAddressTracking) {
@@ -111,13 +111,13 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
this->allocateOrReuseKernelPrivateMemoryIfNeeded(kernel, kernelDescriptor.kernelAttributes.perHwThreadPrivateMemorySize);
if (!launchParams.isIndirect) {
kernel->setGroupCount(threadGroupDimensions->groupCountX,
threadGroupDimensions->groupCountY,
threadGroupDimensions->groupCountZ);
kernel->setGroupCount(threadGroupDimensions.groupCountX,
threadGroupDimensions.groupCountY,
threadGroupDimensions.groupCountZ);
}
if (launchParams.isIndirect && threadGroupDimensions) {
prepareIndirectParams(threadGroupDimensions);
if (launchParams.isIndirect) {
prepareIndirectParams(&threadGroupDimensions);
}
if (kernel->hasIndirectAllocationsAllowed()) {
@@ -182,7 +182,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
kernel, // dispatchInterface
ssh, // surfaceStateHeap
dsh, // dynamicStateHeap
reinterpret_cast<const void *>(threadGroupDimensions), // threadGroupDimensions
reinterpret_cast<const void *>(&threadGroupDimensions), // threadGroupDimensions
nullptr, // outWalkerPtr
&additionalCommands, // additionalCommands
commandListPreemptionMode, // preemptionMode
@@ -300,7 +300,7 @@ inline size_t CommandListCoreFamily<gfxCoreFamily>::estimateBufferSizeMultiTileB
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelSplit(Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
Event *event,
const CmdListKernelLaunchParams &launchParams) {
return appendLaunchKernelWithParams(kernel, threadGroupDimensions, nullptr, launchParams);

18
level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl
View File

@@ -82,7 +82,7 @@ bool CommandListCoreFamily<gfxCoreFamily>::isInOrderNonWalkerSignalingRequired(c
}
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(Kernel *kernel, const ze_group_count_t *threadGroupDimensions, Event *event,
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(Kernel *kernel, const ze_group_count_t &threadGroupDimensions, Event *event,
const CmdListKernelLaunchParams &launchParams) {
if (NEO::DebugManager.flags.ForcePipeControlPriorToWalker.get()) {
@@ -120,7 +120,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
DBG_LOG(PrintDispatchParameters, "Kernel: ", kernelInfo->kernelDescriptor.kernelMetadata.kernelName,
", Group size: ", kernel->getGroupSize()[0], ", ", kernel->getGroupSize()[1], ", ", kernel->getGroupSize()[2],
", Group count: ", threadGroupDimensions->groupCountX, ", ", threadGroupDimensions->groupCountY, ", ", threadGroupDimensions->groupCountZ,
", Group count: ", threadGroupDimensions.groupCountX, ", ", threadGroupDimensions.groupCountY, ", ", threadGroupDimensions.groupCountZ,
", SIMD: ", kernelInfo->getMaxSimdSize());
commandListPerThreadScratchSize = std::max<uint32_t>(commandListPerThreadScratchSize, kernelDescriptor.kernelAttributes.perThreadScratchSize[0]);
@@ -165,13 +165,13 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
kernel->patchGlobalOffset();
this->allocateOrReuseKernelPrivateMemoryIfNeeded(kernel, kernelDescriptor.kernelAttributes.perHwThreadPrivateMemorySize);
if (launchParams.isIndirect && threadGroupDimensions) {
prepareIndirectParams(threadGroupDimensions);
if (launchParams.isIndirect) {
prepareIndirectParams(&threadGroupDimensions);
}
if (!launchParams.isIndirect) {
kernel->setGroupCount(threadGroupDimensions->groupCountX,
threadGroupDimensions->groupCountY,
threadGroupDimensions->groupCountZ);
kernel->setGroupCount(threadGroupDimensions.groupCountX,
threadGroupDimensions.groupCountY,
threadGroupDimensions.groupCountZ);
}
uint64_t eventAddress = 0;
@@ -278,7 +278,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
kernel, // dispatchInterface
ssh, // surfaceStateHeap
dsh, // dynamicStateHeap
reinterpret_cast<const void *>(threadGroupDimensions), // threadGroupDimensions
reinterpret_cast<const void *>(&threadGroupDimensions), // threadGroupDimensions
nullptr, // outWalkerPtr
&additionalCommands, // additionalCommands
kernelPreemptionMode, // preemptionMode
@@ -461,7 +461,7 @@ inline size_t CommandListCoreFamily<gfxCoreFamily>::estimateBufferSizeMultiTileB
template <GFXCORE_FAMILY gfxCoreFamily>
ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelSplit(Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
Event *event,
const CmdListKernelLaunchParams &launchParams) {
if (event) {

4
level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.h
View File

@@ -219,14 +219,14 @@ class AppendFillFixture : public DeviceFixture {
MockCommandList() : WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>() {}
ze_result_t appendLaunchKernelWithParams(Kernel *kernel,
const ze_group_count_t *pThreadGroupDimensions,
const ze_group_count_t &pThreadGroupDimensions,
L0::Event *event,
const CmdListKernelLaunchParams &launchParams) override {
if (numberOfCallsToAppendLaunchKernelWithParams == thresholdOfCallsToAppendLaunchKernelWithParamsToFail) {
return ZE_RESULT_ERROR_UNKNOWN;
}
if (numberOfCallsToAppendLaunchKernelWithParams < 3) {
threadGroupDimensions[numberOfCallsToAppendLaunchKernelWithParams] = *pThreadGroupDimensions;
threadGroupDimensions[numberOfCallsToAppendLaunchKernelWithParams] = pThreadGroupDimensions;
xGroupSizes[numberOfCallsToAppendLaunchKernelWithParams] = kernel->getGroupSize()[0];
}
numberOfCallsToAppendLaunchKernelWithParams++;

62
level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.inl
View File

@@ -137,7 +137,7 @@ void CmdListPipelineSelectStateFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 0;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -180,7 +180,7 @@ void CmdListPipelineSelectStateFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 1;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -200,7 +200,7 @@ void CmdListPipelineSelectStateFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 0;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -247,7 +247,7 @@ void CmdListPipelineSelectStateFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 1;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -267,7 +267,7 @@ void CmdListPipelineSelectStateFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 0;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -290,7 +290,7 @@ void CmdListPipelineSelectStateFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 1;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -336,7 +336,7 @@ void CmdListPipelineSelectStateFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 1;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -398,7 +398,7 @@ void CmdListPipelineSelectStateFixture::testBodyShareStateRegularImmediate() {
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 1;
sizeBefore = regularCommandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = regularCommandListStream.getUsed();
@@ -444,7 +444,7 @@ void CmdListPipelineSelectStateFixture::testBodyShareStateRegularImmediate() {
size_t csrUsedBefore = csrStream.getUsed();
sizeBefore = immediateCmdListStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = immediateCmdListStream.getUsed();
size_t csrUsedAfter = csrStream.getUsed();
@@ -503,7 +503,7 @@ void CmdListPipelineSelectStateFixture::testBodyShareStateImmediateRegular() {
size_t csrUsedBefore = csrStream.getUsed();
sizeBefore = immediateCmdListStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = immediateCmdListStream.getUsed();
size_t csrUsedAfter = csrStream.getUsed();
@@ -544,7 +544,7 @@ void CmdListPipelineSelectStateFixture::testBodyShareStateImmediateRegular() {
auto &cmdQueueStream = commandQueue->commandStream;
sizeBefore = regularCommandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = regularCommandListStream.getUsed();
@@ -587,7 +587,7 @@ void CmdListPipelineSelectStateFixture::testBodySystolicAndScratchOnSecondComman
CmdListKernelLaunchParams launchParams = {};
auto result = ZE_RESULT_SUCCESS;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -596,7 +596,7 @@ void CmdListPipelineSelectStateFixture::testBodySystolicAndScratchOnSecondComman
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.usesSystolicPipelineSelectMode = 1;
mockKernelImmData->kernelDescriptor->kernelAttributes.perThreadScratchSize[0] = 0x40;
result = commandList2->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList2->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList2->close();
@@ -680,7 +680,7 @@ void CmdListThreadArbitrationFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::AgeBased;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -727,7 +727,7 @@ void CmdListThreadArbitrationFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::RoundRobin;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -747,7 +747,7 @@ void CmdListThreadArbitrationFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::AgeBased;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -794,7 +794,7 @@ void CmdListThreadArbitrationFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::RoundRobin;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -814,7 +814,7 @@ void CmdListThreadArbitrationFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::AgeBased;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -837,7 +837,7 @@ void CmdListThreadArbitrationFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::RoundRobinAfterDependency;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -883,7 +883,7 @@ void CmdListThreadArbitrationFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::RoundRobinAfterDependency;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -947,7 +947,7 @@ void CmdListLargeGrfFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.numGrfRequired = GrfConfig::DefaultGrfNumber;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -994,7 +994,7 @@ void CmdListLargeGrfFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.numGrfRequired = GrfConfig::LargeGrfNumber;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -1014,7 +1014,7 @@ void CmdListLargeGrfFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.numGrfRequired = GrfConfig::DefaultGrfNumber;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -1061,7 +1061,7 @@ void CmdListLargeGrfFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.numGrfRequired = GrfConfig::LargeGrfNumber;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -1081,7 +1081,7 @@ void CmdListLargeGrfFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.numGrfRequired = GrfConfig::DefaultGrfNumber;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -1104,7 +1104,7 @@ void CmdListLargeGrfFixture::testBody() {
mockKernelImmData->kernelDescriptor->kernelAttributes.numGrfRequired = GrfConfig::LargeGrfNumber;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -1150,7 +1150,7 @@ void CmdListLargeGrfFixture::testBody() {
{
mockKernelImmData->kernelDescriptor->kernelAttributes.numGrfRequired = GrfConfig::LargeGrfNumber;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();
@@ -1221,12 +1221,12 @@ void ImmediateCmdListSharedHeapsFlushTaskFixtureInit::testBody(NonKernelOperatio
CmdListKernelLaunchParams launchParams = {};
auto result = ZE_RESULT_SUCCESS;
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
validateDispatchFlags(false, ultCsr.recordedImmediateDispatchFlags, ultCsr.recordedSsh);
result = commandListImmediateCoexisting->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediateCoexisting->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
validateDispatchFlags(false, ultCsr.recordedImmediateDispatchFlags, ultCsr.recordedSsh);
@@ -1247,14 +1247,14 @@ void ImmediateCmdListSharedHeapsFlushTaskFixtureInit::testBody(NonKernelOperatio
csrSshHeap->getSpace(csrSshHeap->getAvailableSpace());
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
validateDispatchFlags(false, ultCsr.recordedImmediateDispatchFlags, ultCsr.recordedSsh);
EXPECT_NE(firstSshCpuPointer, sshFirstCmdList->getCpuBase());
result = commandListImmediateCoexisting->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediateCoexisting->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
validateDispatchFlags(false, ultCsr.recordedImmediateDispatchFlags, ultCsr.recordedSsh);

2
level_zero/core/test/unit_tests/gen12lp/test_cmdlist_gen12lp.cpp
View File

@@ -381,7 +381,7 @@ GEN12LPTEST_F(CommandListGen12LpStateComputeModeTrackingTest,
mockKernelImmData->kernelDescriptor->kernelAttributes.threadArbitrationPolicy = NEO::ThreadArbitrationPolicy::RoundRobin;
sizeBefore = commandListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListStream.getUsed();

2
level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp
View File

@@ -28,7 +28,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelWithSLMThenL3IsProgrammedWit
EXPECT_LE(0u, kernel->kernelImmData->getDescriptor().kernelAttributes.slmInlineSize);
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();

8
level_zero/core/test/unit_tests/gen9/test_cmdlist_create_gen9.cpp
View File

@@ -117,7 +117,7 @@ GEN9TEST_F(CommandListCreateGen9, GivenDisabledMidThreadPreemptionWhenLaunchingK
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernelThreadGroup.toHandle(),
&dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(NEO::PreemptionMode::ThreadGroup, commandList->getCommandListPreemptionMode());
auto result = commandList->close();
@@ -147,7 +147,7 @@ GEN9TEST_F(CommandListCreateGen9, GivenUsesFencesForReadWriteImagesWhenLaunching
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernelMidBatch.toHandle(),
&dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(NEO::PreemptionMode::MidBatch, commandList->getCommandListPreemptionMode());
auto result = commandList->close();
@@ -182,11 +182,11 @@ GEN9TEST_F(CommandListCreateGen9, WhenCommandListHasLowerPreemptionLevelThenDoNo
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernelThreadGroup.toHandle(),
&dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(NEO::PreemptionMode::ThreadGroup, commandList->getCommandListPreemptionMode());
commandList->appendLaunchKernel(kernelMidThread.toHandle(),
&dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(NEO::PreemptionMode::ThreadGroup, commandList->getCommandListPreemptionMode());
auto result = commandList->close();

18
level_zero/core/test/unit_tests/mocks/mock_cmdlist.h
View File

@@ -105,7 +105,7 @@ struct WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>
WhiteBox() : ::L0::CommandListCoreFamily<gfxCoreFamily>(BaseClass::defaultNumIddsPerBlock) {}
ze_result_t appendLaunchKernelWithParams(::L0::Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
::L0::Event *event,
const CmdListKernelLaunchParams &launchParams) override {
@@ -118,7 +118,7 @@ struct WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>
ze_result_t appendLaunchMultipleKernelsIndirect(uint32_t numKernels,
const ze_kernel_handle_t *kernelHandles,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
const ze_group_count_t &pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) override {
@@ -128,7 +128,7 @@ struct WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>
}
ze_result_t appendLaunchKernelIndirect(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *pDispatchArgumentsBuffer,
const ze_group_count_t &pDispatchArgumentsBuffer,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch) override {
appendEventKernelIndirectEventHandleValue = hEvent;
@@ -276,21 +276,21 @@ struct MockCommandList : public CommandList {
ADDMETHOD_NOBASE(appendLaunchKernel, ze_result_t, ZE_RESULT_SUCCESS,
(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
ze_event_handle_t hEvent, uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents,
const CmdListKernelLaunchParams &launchParams, bool relaxedOrderingDispatch));
ADDMETHOD_NOBASE(appendLaunchCooperativeKernel, ze_result_t, ZE_RESULT_SUCCESS,
(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *launchKernelArgs,
const ze_group_count_t &launchKernelArgs,
ze_event_handle_t hSignalEvent,
uint32_t numWaitEvents,
ze_event_handle_t *waitEventHandles, bool relaxedOrderingDispatch));
ADDMETHOD_NOBASE(appendLaunchKernelIndirect, ze_result_t, ZE_RESULT_SUCCESS,
(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *pDispatchArgumentsBuffer,
const ze_group_count_t &DispatchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch));
@@ -299,7 +299,7 @@ struct MockCommandList : public CommandList {
(uint32_t numKernels,
const ze_kernel_handle_t *kernelHandles,
const uint32_t *pNumLaunchArguments,
const ze_group_count_t *pLaunchArgumentsBuffer,
const ze_group_count_t &pLaunchArgumentsBuffer,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents, bool relaxedOrderingDispatch));
@@ -672,7 +672,7 @@ class MockCommandListForAppendLaunchKernel : public WhiteBox<::L0::CommandListCo
public:
CmdListHelper cmdListHelper;
ze_result_t appendLaunchKernel(ze_kernel_handle_t kernelHandle,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
ze_event_handle_t hEvent,
uint32_t numWaitEvents,
ze_event_handle_t *phWaitEvents,
@@ -682,7 +682,7 @@ class MockCommandListForAppendLaunchKernel : public WhiteBox<::L0::CommandListCo
cmdListHelper.isaAllocation = kernel->getIsaAllocation();
cmdListHelper.residencyContainer = kernel->getResidencyContainer();
cmdListHelper.groupSize = kernel->getGroupSize();
cmdListHelper.threadGroupDimensions = *threadGroupDimensions;
cmdListHelper.threadGroupDimensions = threadGroupDimensions;
auto kernelName = kernel->getImmutableData()->getDescriptor().kernelMetadata.kernelName;
NEO::ArgDescriptor arg;

2
level_zero/core/test/unit_tests/os_interface/linux/cmdqueue_linux_tests.cpp
View File

@@ -63,7 +63,7 @@ HWTEST2_F(CommandQueueLinuxTests, givenExecBufferErrorOnXeHpcWhenExecutingComman
ze_group_count_t dispatchFunctionArguments{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel.toHandle(), &dispatchFunctionArguments, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), dispatchFunctionArguments, nullptr, 0, nullptr, launchParams, false);
commandList->close();
ze_command_list_handle_t cmdListHandles[1] = {commandList->toHandle()};

12
level_zero/core/test/unit_tests/sources/assert/test_assert.cpp
View File

@@ -155,7 +155,7 @@ TEST(CommandListAssertTest, GivenCmdListWhenKernelWithAssertAppendedThenHasKerne
kernel.descriptor.kernelAttributes.flags.usesAssert = true;
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_TRUE(commandList->hasKernelWithAssert());
@@ -197,7 +197,7 @@ TEST_F(CommandListImmediateWithAssert, GivenImmediateCmdListWhenKernelWithAssert
kernel.descriptor.kernelAttributes.flags.usesAssert = true;
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, assertHandler->printAssertAndAbortCalled);
@@ -263,7 +263,7 @@ HWTEST2_F(CommandListImmediateWithAssert, givenKernelWithAssertWhenAppendedToAsy
kernel.descriptor.kernelAttributes.flags.usesAssert = true;
CmdListKernelLaunchParams launchParams = {};
result = cmdList.appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = cmdList.appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(0u, cmdList.checkAssertCalled);
@@ -296,7 +296,7 @@ HWTEST2_F(CommandListImmediateWithAssert, givenKernelWithAssertWhenAppendedToSyn
kernel.descriptor.kernelAttributes.flags.usesAssert = true;
CmdListKernelLaunchParams launchParams = {};
result = cmdList.appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = cmdList.appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, cmdList.checkAssertCalled);
@@ -333,7 +333,7 @@ HWTEST2_F(CommandListImmediateWithAssert, givenKernelWithAssertWhenAppendToSynch
kernel.descriptor.kernelAttributes.flags.usesAssert = true;
CmdListKernelLaunchParams launchParams = {};
result = cmdList.appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = cmdList.appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_DEVICE_LOST, result);
EXPECT_EQ(1u, cmdList.checkAssertCalled);
@@ -365,7 +365,7 @@ TEST_F(CommandQueueWithAssert, GivenCmdListWithAssertWhenExecutingThenCommandQue
kernel.descriptor.kernelAttributes.flags.usesAssert = true;
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
commandList->close();

42
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp
View File

@@ -1097,12 +1097,12 @@ HWTEST2_F(CommandListCreate, givenDirectSubmissionAndImmCmdListWhenDispatchingTh
EXPECT_EQ(ultCsr->latestFlushedBatchBuffer.hasStallingCmds, bbFlag);
};
// non-pipelined state
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false), false, true);
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false), false, true);
// non-pipelined state already programmed
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false), false, false);
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false), false, false);
verifyFlags(commandList->appendLaunchKernelIndirect(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, false), false, false);
verifyFlags(commandList->appendLaunchKernelIndirect(kernel.toHandle(), groupCount, nullptr, 0, nullptr, false), false, false);
verifyFlags(commandList->appendBarrier(nullptr, 0, nullptr, false), true, true);
@@ -1144,7 +1144,7 @@ HWTEST2_F(CommandListCreate, givenDirectSubmissionAndImmCmdListWhenDispatchingTh
const void **ranges = reinterpret_cast<const void **>(&dstPtr[0]);
verifyFlags(commandList->appendMemoryRangesBarrier(1, &rangeSizes, ranges, nullptr, 0, nullptr), true, true);
verifyFlags(commandList->appendLaunchCooperativeKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, false), false, false);
verifyFlags(commandList->appendLaunchCooperativeKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, false), false, false);
driverHandle->releaseImportedPointer(dstPtr);
}
@@ -1223,14 +1223,14 @@ HWTEST2_F(CommandListCreate, givenDirectSubmissionAndImmCmdListWhenDispatchingDi
// non-pipelined state or first in-order exec
resetFlags();
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 1, &event, launchParams, false));
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 1, &event, launchParams, false));
// non-pipelined state already programmed
resetFlags();
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, numWaitEvents, waitlist, launchParams, false));
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, numWaitEvents, waitlist, launchParams, false));
resetFlags();
verifyFlags(commandList->appendLaunchKernelIndirect(kernel.toHandle(), &groupCount, nullptr, numWaitEvents, waitlist, false));
verifyFlags(commandList->appendLaunchKernelIndirect(kernel.toHandle(), groupCount, nullptr, numWaitEvents, waitlist, false));
resetFlags();
verifyFlags(commandList->appendMemoryCopy(dstPtr, srcPtr, 8, nullptr, numWaitEvents, waitlist, false, false));
@@ -1263,7 +1263,7 @@ HWTEST2_F(CommandListCreate, givenDirectSubmissionAndImmCmdListWhenDispatchingDi
}
resetFlags();
verifyFlags(commandList->appendLaunchCooperativeKernel(kernel.toHandle(), &groupCount, nullptr, numWaitEvents, waitlist, false));
verifyFlags(commandList->appendLaunchCooperativeKernel(kernel.toHandle(), groupCount, nullptr, numWaitEvents, waitlist, false));
}
driverHandle->releaseImportedPointer(dstPtr);
@@ -1288,7 +1288,7 @@ HWTEST2_F(CommandListCreate, whenDispatchingThenPassNumCsrClients, IsAtLeastXeHp
ultCsr->registerClient(&client1);
ultCsr->registerClient(&client2);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(ultCsr->latestFlushedBatchBuffer.numCsrClients, ultCsr->getNumClients());
@@ -1326,7 +1326,7 @@ HWTEST_F(CommandListCreate, givenSignalEventWhenCallingSynchronizeThenUnregister
EXPECT_EQ(ultCsr->getNumClients(), 0u);
{
commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event1, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event1, 0, nullptr, launchParams, false);
EXPECT_EQ(ultCsr->getNumClients(), 1u);
Event::fromHandle(event1)->setIsCompleted();
@@ -1336,7 +1336,7 @@ HWTEST_F(CommandListCreate, givenSignalEventWhenCallingSynchronizeThenUnregister
}
{
commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event2, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event2, 0, nullptr, launchParams, false);
EXPECT_EQ(ultCsr->getNumClients(), 1u);
*reinterpret_cast<uint32_t *>(Event::fromHandle(event2)->getHostAddress()) = static_cast<uint32_t>(Event::STATE_SIGNALED);
@@ -1346,7 +1346,7 @@ HWTEST_F(CommandListCreate, givenSignalEventWhenCallingSynchronizeThenUnregister
}
{
commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event3, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event3, 0, nullptr, launchParams, false);
EXPECT_EQ(ultCsr->getNumClients(), 1u);
zeEventHostReset(event3);
@@ -1389,7 +1389,7 @@ HWTEST_F(CommandListCreate, givenDebugFlagSetWhenCallingSynchronizeThenDontUnreg
ASSERT_EQ(ZE_RESULT_SUCCESS, eventPool->createEvent(&eventDesc, &event));
EXPECT_EQ(ultCsr->getNumClients(), 0u);
commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event, 0, nullptr, launchParams, false);
EXPECT_EQ(ultCsr->getNumClients(), 1u);
Event::fromHandle(event)->setIsCompleted();
@@ -1462,10 +1462,10 @@ HWTEST2_F(CommandListCreate, givenDirectSubmissionAndImmCmdListWhenDispatchingTh
ze_event_handle_t *waitlist = hasEventDependencies ? &event : nullptr;
uint32_t numWaitlistEvents = hasEventDependencies ? 1 : 0;
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, numWaitlistEvents, waitlist, launchParams, false),
verifyFlags(commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, numWaitlistEvents, waitlist, launchParams, false),
hasEventDependencies, hasEventDependencies);
verifyFlags(commandList->appendLaunchKernelIndirect(kernel.toHandle(), &groupCount, nullptr, numWaitlistEvents, waitlist, false),
verifyFlags(commandList->appendLaunchKernelIndirect(kernel.toHandle(), groupCount, nullptr, numWaitlistEvents, waitlist, false),
hasEventDependencies, hasEventDependencies);
verifyFlags(commandList->appendBarrier(nullptr, numWaitlistEvents, waitlist, false),
@@ -1522,7 +1522,7 @@ HWTEST2_F(CommandListCreate, givenDirectSubmissionAndImmCmdListWhenDispatchingTh
for (bool hasEventDependencies : {true, false}) {
ze_event_handle_t *waitlist = hasEventDependencies ? &event : nullptr;
uint32_t numWaitlistEvents = hasEventDependencies ? 1 : 0;
verifyFlags(commandList->appendLaunchCooperativeKernel(kernel.toHandle(), &groupCount, nullptr, numWaitlistEvents, waitlist, false),
verifyFlags(commandList->appendLaunchCooperativeKernel(kernel.toHandle(), groupCount, nullptr, numWaitlistEvents, waitlist, false),
hasEventDependencies, hasEventDependencies);
}
@@ -1571,9 +1571,9 @@ HWTEST2_F(CommandListCreate, givenInOrderExecutionWhenDispatchingRelaxedOrdering
ultCsr->registerClient(&client1);
ultCsr->registerClient(&client2);
commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
if (useImmediateFlushTask) {
EXPECT_TRUE(ultCsr->recordedImmediateDispatchFlags.hasRelaxedOrderingDependencies);
} else {
@@ -1750,12 +1750,12 @@ HWTEST2_F(CommandListCreate, givenInOrderExecutionWhenDispatchingRelaxedOrdering
auto cmdStream = cmdList->getCmdContainer().getCommandStream();
cmdList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
cmdList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
cmdList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
cmdList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t offset = cmdStream->getUsed();
cmdList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
cmdList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
GenCmdList genCmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(

38
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp
View File

@@ -1446,7 +1446,7 @@ HWTEST_F(PrimaryBatchBufferCmdListTest, givenPrimaryBatchBufferWhenAppendingKern
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t cmdListUsed = cmdListStream.getUsed();
@@ -1475,9 +1475,9 @@ HWTEST_F(PrimaryBatchBufferCmdListTest, givenPrimaryBatchBufferWhenAppendingKern
result = commandList->reset();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
cmdListUsed = cmdListStream.getUsed();
@@ -1500,7 +1500,7 @@ HWTEST_F(PrimaryBatchBufferCmdListTest, givenPrimaryBatchBufferWhenCommandListHa
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto firstChainBufferAllocation = cmdListStream.getGraphicsAllocation();
@@ -1508,9 +1508,9 @@ HWTEST_F(PrimaryBatchBufferCmdListTest, givenPrimaryBatchBufferWhenCommandListHa
size_t firstCmdBufferUsed = cmdListStream.getUsed();
auto bbStartSpace = ptrOffset(cmdListStream.getCpuBase(), firstCmdBufferUsed);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto secondChainBufferAllocation = cmdListStream.getGraphicsAllocation();
@@ -1540,7 +1540,7 @@ HWTEST_F(PrimaryBatchBufferCmdListTest, givenRegularCmdListWhenFlushingThenPassS
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
EXPECT_EQ(ZE_RESULT_SUCCESS, commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false));
EXPECT_EQ(ZE_RESULT_SUCCESS, commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false));
EXPECT_EQ(ZE_RESULT_SUCCESS, commandList->close());
@@ -1553,7 +1553,7 @@ HWTEST_F(PrimaryBatchBufferCmdListTest, givenRegularCmdListWhenFlushingThenPassS
HWTEST_F(PrimaryBatchBufferCmdListTest, givenCmdListWhenCallingSynchronizeThenUnregisterCsrClient) {
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
EXPECT_EQ(ZE_RESULT_SUCCESS, commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false));
EXPECT_EQ(ZE_RESULT_SUCCESS, commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false));
EXPECT_EQ(ZE_RESULT_SUCCESS, commandList->close());
@@ -1689,24 +1689,24 @@ HWTEST2_F(PrimaryBatchBufferPreamblelessCmdListTest,
// command list 1 will have two kernels, transition from cached MOCS to uncached MOCS state
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
kernel->kernelRequiresUncachedMocsCount++;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
// command list 2 will have two kernels, transition from uncached MOCS to cached MOCS state
result = commandList2->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList2->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
kernel->kernelRequiresUncachedMocsCount--;
result = commandList2->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList2->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList2->close();
@@ -1750,7 +1750,7 @@ HWTEST2_F(PrimaryBatchBufferPreamblelessCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -1789,7 +1789,7 @@ HWTEST2_F(PrimaryBatchBufferPreamblelessCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -1808,13 +1808,13 @@ HWTEST2_F(PrimaryBatchBufferPreamblelessCmdListTest,
size_t queueUsedSize = cmdQueueStream.getUsed();
auto gpuReturnAddress = cmdQueueStream.getGpuBase() + queueUsedSize;
result = commandList2->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList2->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList2->close();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList3->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList3->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList3->close();
@@ -1865,7 +1865,7 @@ HWTEST2_F(PrimaryBatchBufferPreamblelessCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -1886,13 +1886,13 @@ HWTEST2_F(PrimaryBatchBufferPreamblelessCmdListTest,
kernel->kernelRequiresUncachedMocsCount++;
result = commandList2->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList2->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList2->close();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList3->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList3->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList3->close();

20
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_3.cpp
View File

@@ -1537,7 +1537,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &commandsToPatch = commandList->commandsToPatch;
@@ -1547,7 +1547,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 1;
size_t usedBefore = cmdStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.disableEuFusion) {
@@ -1566,7 +1566,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
EXPECT_EQ(0u, commandsToPatch.size());
}
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.disableEuFusion) {
@@ -1578,7 +1578,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 0;
usedBefore = cmdStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.disableEuFusion) {
@@ -1598,7 +1598,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 1;
usedBefore = cmdStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.disableEuFusion) {
@@ -1665,13 +1665,13 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &commandsToPatch = commandList->commandsToPatch;
EXPECT_EQ(0u, commandsToPatch.size());
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.computeDispatchAllWalker) {
@@ -1687,7 +1687,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
EXPECT_EQ(0u, commandsToPatch.size());
}
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.computeDispatchAllWalker) {
@@ -1696,7 +1696,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
EXPECT_EQ(0u, commandsToPatch.size());
}
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.computeDispatchAllWalker) {
@@ -1712,7 +1712,7 @@ HWTEST2_F(FrontEndPrimaryBatchBufferCommandListTest,
EXPECT_EQ(0u, commandsToPatch.size());
}
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.computeDispatchAllWalker) {

8
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp
View File

@@ -331,7 +331,7 @@ HWTEST_F(CommandListImmediateFlushTaskComputeTests, givenUseCsrImmediateSubmissi
std::unique_ptr<L0::CommandList> commandList(CommandList::createImmediate(productFamily, device, &queueDesc, false, NEO::EngineGroupType::Compute, returnValue));
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
}
@@ -367,13 +367,13 @@ HWTEST2_F(CommandListAppendLaunchKernelResetKernelCount, givenIsKernelSplitOpera
event->increaseKernelCount();
launchParams.isKernelSplitOperation = true;
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(2u, event->getKernelCount());
}
{
launchParams.isKernelSplitOperation = false;
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, event->getKernelCount());
}
@@ -1451,7 +1451,7 @@ HWTEST2_F(ImmediateCmdListSharedHeapsImmediateFlushTaskTest,
auto result = ZE_RESULT_SUCCESS;
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
csrUsedAfter = csrStream.getUsed();

70
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_5.cpp
View File

@@ -1121,7 +1121,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
@@ -1288,7 +1288,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
@@ -1341,7 +1341,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
container.getHeapWithRequiredSizeAndAlignment(NEO::HeapType::DYNAMIC_STATE, dshHeap->getMaxAvailableSpace(), 0);
}
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
ssBaseAddress = sshHeap->getGpuBase();
@@ -1495,7 +1495,7 @@ HWTEST2_F(CommandListBindlessSshPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
@@ -1531,7 +1531,7 @@ HWTEST2_F(CommandListBindlessSshPrivateHeapTest,
size_t cmdListBefore = cmdListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
ssBaseAddress = sshHeap->getGpuBase();
@@ -1643,11 +1643,11 @@ HWTEST2_F(CommandListBindlessSshPrivateHeapTest,
CmdListKernelLaunchParams launchParams = {};
memset(kernel->dynamicStateHeapData.get(), 0, kernel->dynamicStateHeapDataSize);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -1716,7 +1716,7 @@ HWTEST2_F(CommandListBindlessSshPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(mockKernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(mockKernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
@@ -1796,7 +1796,7 @@ HWTEST2_F(CommandListBindlessSshPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(mockKernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(mockKernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
@@ -1851,7 +1851,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
@@ -1901,7 +1901,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
container.getHeapWithRequiredSizeAndAlignment(NEO::HeapType::DYNAMIC_STATE, dshHeap->getMaxAvailableSpace(), 0);
}
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
ssBaseAddress = sshHeap->getGpuBase();
@@ -1947,7 +1947,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
EXPECT_EQ((statlessMocs << 1), sbaCmd->getStatelessDataPortAccessMemoryObjectControlState());
size_t sizeBefore = cmdListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
cmdList.clear();
@@ -1982,7 +1982,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
size_t cmdListUsedBefore = 0;
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t cmdListUsedAfter = cmdListImmediateStream.getUsed();
size_t csrUsedAfter = csrStream.getUsed();
@@ -2077,7 +2077,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
}
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
csrUsedAfter = csrStream.getUsed();
@@ -2134,7 +2134,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList cmdList;
@@ -2262,7 +2262,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
auto &csrStream = csrImmediate.commandStream;
size_t csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -2350,7 +2350,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
auto &csrStream = csrImmediate.commandStream;
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -2416,7 +2416,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
auto &container = commandList->getCmdContainer();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto sshHeap = container.getIndirectHeap(NEO::HeapType::SURFACE_STATE);
@@ -2540,7 +2540,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t usedBefore = cmdStream.getUsed();
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t usedAfter = cmdStream.getUsed();
@@ -2590,7 +2590,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
auto &csrStream = csrImmediate.commandStream;
size_t csrBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrAfter = csrStream.getUsed();
@@ -2618,7 +2618,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
uint32_t uncachedStatlessMocs = getMocs(false);
@@ -2661,7 +2661,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
kernel->kernelRequiresUncachedMocsCount = 1;
size_t csrBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrAfter = csrStream.getUsed();
@@ -2703,7 +2703,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
size_t csrBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrAfter = csrStream.getUsed();
@@ -2723,7 +2723,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
kernel->kernelRequiresUncachedMocsCount = 0;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(static_cast<int32_t>(cachedStatlessMocs), requiredState.statelessMocs.value);
@@ -2777,7 +2777,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
size_t csrBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrAfter = csrStream.getUsed();
@@ -2797,7 +2797,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
kernel->kernelRequiresUncachedMocsCount = 1;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(static_cast<int32_t>(uncachedStatlessMocs), requiredState.statelessMocs.value);
@@ -2886,7 +2886,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -2920,7 +2920,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -2949,7 +2949,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -2985,7 +2985,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
result = commandList->reset();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -3019,7 +3019,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -3058,10 +3058,10 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
auto &cmdListStream = *container.getCommandStream();
size_t usedBefore = cmdListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t usedAfter = cmdListStream.getUsed();
@@ -3108,7 +3108,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
cmdListObject->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
cmdListObject->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = cmdListObject->close();
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
@@ -3136,7 +3136,7 @@ HWTEST2_F(CommandListStateBaseAddressPrivateHeapTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
cmdListObject->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
cmdListObject->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = cmdListObject->close();
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);

42
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_6.cpp
View File

@@ -1079,7 +1079,7 @@ HWTEST2_F(ImmediateCmdListSharedHeapsTest, givenMultipleCommandListsUsingSharedH
size_t sshUsed = csrSshHeap->getUsed();
size_t csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -1136,7 +1136,7 @@ HWTEST2_F(ImmediateCmdListSharedHeapsTest, givenMultipleCommandListsUsingSharedH
sshUsed = csrSshHeap->getUsed();
csrUsedBefore = csrStream.getUsed();
result = commandListImmediateCoexisting->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediateCoexisting->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
csrUsedAfter = csrStream.getUsed();
@@ -1213,7 +1213,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &container = commandList->getCmdContainer();
@@ -1326,7 +1326,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
size_t csrUsedBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -1387,7 +1387,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &container = commandList->getCmdContainer();
@@ -1442,7 +1442,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
auto &csrStream = csrImmediate.commandStream;
size_t csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -1467,7 +1467,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
size_t csrUsedBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -1505,7 +1505,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
EXPECT_EQ((statlessMocs << 1), sbaCmd->getStatelessDataPortAccessMemoryObjectControlState());
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -1535,7 +1535,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &container = commandList->getCmdContainer();
@@ -1633,7 +1633,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
EXPECT_EQ((statlessMocs << 1), sbaCmd->getStatelessDataPortAccessMemoryObjectControlState());
result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &containerPrivateHeap = commandListPrivateHeap->getCmdContainer();
@@ -1756,7 +1756,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &containerPrivateHeap = commandListPrivateHeap->getCmdContainer();
@@ -1879,7 +1879,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
EXPECT_EQ((statlessMocs << 1), sbaCmd->getStatelessDataPortAccessMemoryObjectControlState());
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &requiredState = commandList->requiredStreamState.stateBaseAddress;
@@ -1992,7 +1992,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
size_t csrUsedBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -2045,7 +2045,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
EXPECT_EQ((statlessMocs << 1), sbaCmd->getStatelessDataPortAccessMemoryObjectControlState());
result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &containerPrivateHeap = commandListPrivateHeap->getCmdContainer();
@@ -2170,7 +2170,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListPrivateHeap->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto &containerPrivateHeap = commandListPrivateHeap->getCmdContainer();
@@ -2298,7 +2298,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
auto globalSurfaceHeap = csrImmediate.getGlobalStatelessHeap();
size_t csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -2361,7 +2361,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -2427,7 +2427,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
size_t csrUsedBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrUsedAfter = csrStream.getUsed();
@@ -2476,7 +2476,7 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
cmdListObject->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
cmdListObject->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = cmdListObject->close();
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
@@ -2513,14 +2513,14 @@ HWTEST2_F(CommandListStateBaseAddressGlobalStatelessTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
ptrArg.as<ArgDescPointer>().bindless = undefined<CrossThreadDataOffset>;
ptrArg.as<ArgDescPointer>().bindful = 0x40;
mockKernelImmData->kernelDescriptor->payloadMappings.explicitArgs[0] = ptrArg;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
}

81
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_7.cpp
View File

@@ -264,7 +264,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenSignalEventWhenAppendLaunchCoopera
auto commandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandList->initialize(device, NEO::EngineGroupType::RenderCompute, 0u);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, false);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
EXPECT_EQ(event.get(), commandList->appendKernelEventValue);
}
@@ -289,13 +289,14 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenSignalEventWhenAppendLaunchMultipl
commandList->initialize(device, NEO::EngineGroupType::RenderCompute, 0u);
const ze_kernel_handle_t launchKernels = kernel->toHandle();
const ze_group_count_t launchKernelArgs = {1, 1, 1};
uint32_t *numLaunchArgs;
ze_device_mem_alloc_desc_t deviceDesc = {};
returnValue = context->allocDeviceMem(
device->toHandle(), &deviceDesc, 16384u, 4096u, reinterpret_cast<void **>(&numLaunchArgs));
ASSERT_EQ(ZE_RESULT_SUCCESS, returnValue);
returnValue = commandList->appendLaunchMultipleKernelsIndirect(1, &launchKernels, numLaunchArgs, nullptr, event->toHandle(), 0, nullptr, false);
returnValue = commandList->appendLaunchMultipleKernelsIndirect(1, &launchKernels, numLaunchArgs, launchKernelArgs, event->toHandle(), 0, nullptr, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, returnValue);
EXPECT_EQ(event->toHandle(), commandList->appendEventMultipleKernelIndirectEventHandleValue);
@@ -330,7 +331,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenSignalEventWhenAppendLaunchIndirec
auto result = context->allocDeviceMem(device->toHandle(), &deviceDesc, 16384u, 4096u, &alloc);
ASSERT_EQ(result, ZE_RESULT_SUCCESS);
result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), static_cast<ze_group_count_t *>(alloc), event->toHandle(), 0, nullptr, false);
result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), *static_cast<ze_group_count_t *>(alloc), event->toHandle(), 0, nullptr, false);
EXPECT_EQ(result, ZE_RESULT_SUCCESS);
EXPECT_EQ(event->toHandle(), commandList->appendEventKernelIndirectEventHandleValue);
@@ -351,7 +352,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, GivenComputeModePropertiesWhenUpdateStr
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x100;
const ze_group_count_t launchKernelArgs = {};
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
if (commandList->stateComputeModeTracking) {
if (productHelper.getScmPropertyCoherencyRequiredSupport()) {
EXPECT_EQ(0, commandList->finalStreamState.stateComputeMode.isCoherencyRequired.value);
@@ -369,7 +370,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, GivenComputeModePropertiesWhenUpdateStr
}
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x80;
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
if constexpr (TestTraits<gfxCoreFamily>::largeGrfModeInStateComputeModeSupported) {
EXPECT_EQ(productHelper.isGrfNumReportedWithScm(), commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
}
@@ -403,7 +404,7 @@ HWTEST2_F(CommandListAppendLaunchKernel,
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x100;
const ze_group_count_t launchKernelArgs = {};
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
if (commandList->stateComputeModeTracking) {
if (productHelper.getScmPropertyCoherencyRequiredSupport()) {
EXPECT_EQ(0, commandList->finalStreamState.stateComputeMode.isCoherencyRequired.value);
@@ -425,7 +426,7 @@ HWTEST2_F(CommandListAppendLaunchKernel,
}
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x80;
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(productHelper.isGrfNumReportedWithScm(), commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
if (productHelper.getScmPropertyCoherencyRequiredSupport()) {
EXPECT_EQ(0, commandList->finalStreamState.stateComputeMode.isCoherencyRequired.value);
@@ -448,7 +449,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, GivenComputeModePropertiesWhenPropertes
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x100;
const ze_group_count_t launchKernelArgs = {};
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
if (commandList->stateComputeModeTracking) {
if (productHelper.getScmPropertyCoherencyRequiredSupport()) {
EXPECT_EQ(0, commandList->finalStreamState.stateComputeMode.isCoherencyRequired.value);
@@ -465,7 +466,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, GivenComputeModePropertiesWhenPropertes
EXPECT_EQ(productHelper.isGrfNumReportedWithScm(), commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
}
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
if (productHelper.getScmPropertyCoherencyRequiredSupport()) {
EXPECT_EQ(0, commandList->finalStreamState.stateComputeMode.isCoherencyRequired.value);
} else {
@@ -618,7 +619,7 @@ HWTEST2_F(CmdlistAppendLaunchKernelTests,
ze_group_count_t groupCount = {3, 2, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(scratchPerThreadSize, commandList->getCommandListPerThreadScratchSize());
@@ -667,7 +668,7 @@ HWTEST2_F(CmdlistAppendLaunchKernelTests,
ze_group_count_t groupCount = {3, 2, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(scratchPerThreadSize, commandList->getCommandListPerThreadScratchSize());
@@ -737,7 +738,7 @@ HWTEST2_F(CmdlistAppendLaunchKernelTests,
ze_event_handle_t eventHandles[1] = {event->toHandle()};
EXPECT_EQ(MockEvent::STATE_CLEARED, static_cast<MockEvent *>(event.get())->isCompleted);
result = CommandList::fromHandle(cmdListHandle)->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, eventHandles, launchParams, false);
result = CommandList::fromHandle(cmdListHandle)->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, eventHandles, launchParams, false);
EXPECT_EQ(result, ZE_RESULT_SUCCESS);
EXPECT_EQ(MockEvent::STATE_SIGNALED, static_cast<MockEvent *>(event.get())->isCompleted);
@@ -760,12 +761,12 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 1;
size_t usedBefore = cmdStream.getUsed();
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t usedAfter = cmdStream.getUsed();
GenCmdList cmdList;
@@ -798,7 +799,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
}
usedBefore = cmdStream.getUsed();
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
usedAfter = cmdStream.getUsed();
cmdList.clear();
@@ -826,7 +827,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
cmdStream.getSpace(cmdStream.getAvailableSpace() - sizeof(MI_BATCH_BUFFER_END));
auto oldCmdBuffer = cmdStream.getGraphicsAllocation();
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
usedBefore = 0;
usedAfter = cmdStream.getUsed();
@@ -866,7 +867,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
void *oldBase = cmdStream.getCpuBase();
oldCmdBuffer = cmdStream.getGraphicsAllocation();
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
newCmdBuffer = cmdStream.getGraphicsAllocation();
ASSERT_NE(oldCmdBuffer, newCmdBuffer);
@@ -929,10 +930,10 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t usedBefore = cmdStream.getUsed();
commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
size_t usedAfter = cmdStream.getUsed();
GenCmdList cmdList;
@@ -962,7 +963,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
}
usedBefore = cmdStream.getUsed();
commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
usedAfter = cmdStream.getUsed();
cmdList.clear();
@@ -988,7 +989,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
void *oldBase = cmdStream.getCpuBase();
cmdStream.getSpace(cmdStream.getAvailableSpace() - 2 * sizeof(MI_BATCH_BUFFER_END));
usedBefore = cmdStream.getUsed();
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto newCmdBuffer = cmdStream.getGraphicsAllocation();
ASSERT_NE(oldCmdBuffer, newCmdBuffer);
@@ -1034,7 +1035,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenFrontEndTrackingIsUsedWhenPro
oldCmdBuffer = cmdStream.getGraphicsAllocation();
usedBefore = 0;
commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
usedAfter = cmdStream.getUsed();
newCmdBuffer = cmdStream.getGraphicsAllocation();
@@ -1094,27 +1095,27 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 1;
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 0;
cmdListStream.getSpace(cmdListStream.getAvailableSpace() - sizeof(MI_BATCH_BUFFER_END));
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 1;
cmdListStream.getSpace(cmdListStream.getAvailableSpace() - 2 * sizeof(MI_BATCH_BUFFER_END));
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.disableEuFusion) {
@@ -1347,23 +1348,23 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
cmdListStream.getSpace(cmdListStream.getAvailableSpace() - 2 * sizeof(MI_BATCH_BUFFER_END));
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
cmdListStream.getSpace(cmdListStream.getAvailableSpace() - sizeof(MI_BATCH_BUFFER_END));
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
result = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
if (fePropertiesSupport.computeDispatchAllWalker) {
@@ -1593,7 +1594,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenCmdQueueAndImmediateCmdListUs
mockKernelImmData->kernelDescriptor->kernelAttributes.flags.requiresDisabledEUFusion = 1;
size_t usedBefore = regularCmdListStream.getUsed();
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t usedAfter = regularCmdListStream.getUsed();
@@ -1657,7 +1658,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenCmdQueueAndImmediateCmdListUs
size_t csrUsedBefore = csrStream.getUsed();
usedBefore = immediateCmdListStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
usedAfter = immediateCmdListStream.getUsed();
size_t csrUsedAfter = csrStream.getUsed();
@@ -1725,7 +1726,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenCmdQueueAndImmediateCmdListUs
size_t csrUsedBefore = csrStream.getUsed();
size_t usedBefore = immediateCmdListStream.getUsed();
ze_result_t result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
ze_result_t result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t usedAfter = immediateCmdListStream.getUsed();
size_t csrUsedAfter = csrStream.getUsed();
@@ -1771,7 +1772,7 @@ HWTEST2_F(FrontEndMultiReturnCommandListTest, givenCmdQueueAndImmediateCmdListUs
auto &regularCmdListStream = *commandList->getCmdContainer().getCommandStream();
usedBefore = regularCmdListStream.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
usedAfter = regularCmdListStream.getUsed();
@@ -2845,7 +2846,7 @@ template <GFXCORE_FAMILY gfxCoreFamily>
class MockCommandListHwKernelSplit : public WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>> {
public:
ze_result_t appendLaunchKernelSplit(::L0::Kernel *kernel,
const ze_group_count_t *threadGroupDimensions,
const ze_group_count_t &threadGroupDimensions,
::L0::Event *event,
const CmdListKernelLaunchParams &launchParams) override {
passedKernel = kernel;
@@ -3086,7 +3087,7 @@ HWTEST2_F(CommandListMappedTimestampTest, givenMappedTimestampSignalEventWhenApp
auto commandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandList->initialize(device, NEO::EngineGroupType::RenderCompute, 0u);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, false);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
EXPECT_EQ(event.get(), commandList->peekMappedEventList()[0]);
}
@@ -3112,7 +3113,7 @@ HWTEST2_F(CommandListMappedTimestampTest, givenSignalEventWithoutMappedTimstampW
auto commandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandList->initialize(device, NEO::EngineGroupType::RenderCompute, 0u);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, false);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
EXPECT_EQ(0u, commandList->peekMappedEventList().size());
}
@@ -3138,7 +3139,7 @@ HWTEST2_F(CommandListMappedTimestampTest, givenMappedTimestampSignalEventWhenApp
auto commandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandList->initialize(device, NEO::EngineGroupType::RenderCompute, 0u);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, false);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
returnValue = commandList->appendBarrier(event->toHandle(), 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);

66
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_1.cpp
View File

@@ -43,7 +43,7 @@ HWTEST_F(CommandListAppendLaunchKernelMockModule, givenKernelWithIndirectAllocat
ze_result_t returnValue;
CmdListKernelLaunchParams launchParams = {};
{
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, returnValue);
EXPECT_TRUE(commandList->hasIndirectAllocationsAllowed());
}
@@ -55,7 +55,7 @@ HWTEST_F(CommandListAppendLaunchKernelMockModule, givenKernelWithIndirectAllocat
kernel->unifiedMemoryControls.indirectSharedAllocationsAllowed = true;
kernel->unifiedMemoryControls.indirectHostAllocationsAllowed = false;
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, returnValue);
EXPECT_TRUE(commandList->hasIndirectAllocationsAllowed());
}
@@ -67,7 +67,7 @@ HWTEST_F(CommandListAppendLaunchKernelMockModule, givenKernelWithIndirectAllocat
kernel->unifiedMemoryControls.indirectSharedAllocationsAllowed = false;
kernel->unifiedMemoryControls.indirectHostAllocationsAllowed = false;
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, returnValue);
EXPECT_TRUE(commandList->hasIndirectAllocationsAllowed());
}
@@ -84,7 +84,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithIndirectAllocationsNotAll
ze_result_t returnValue;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
ASSERT_FALSE(commandList->hasIndirectAllocationsAllowed());
@@ -128,7 +128,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithThreadArbitrationPolicySe
ze_result_t returnValue;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
ASSERT_EQ(NEO::ThreadArbitrationPolicy::RoundRobin, commandList->getFinalStreamState().stateComputeMode.threadArbitrationPolicy.value);
@@ -150,7 +150,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithThreadArbitrationPolicySe
ze_result_t returnValue;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
ASSERT_EQ(NEO::ThreadArbitrationPolicy::AgeBased, commandList->getFinalStreamState().stateComputeMode.threadArbitrationPolicy.value);
@@ -227,7 +227,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithPrintfUsedWhenAppendedToC
EXPECT_TRUE(kernel->kernelImmData->getDescriptor().kernelAttributes.flags.usesPrintf);
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, commandList->getPrintfKernelContainer().size());
@@ -242,13 +242,13 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithPrintfUsedWhenAppendedToC
EXPECT_TRUE(kernel->kernelImmData->getDescriptor().kernelAttributes.flags.usesPrintf);
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, commandList->getPrintfKernelContainer().size());
EXPECT_EQ(kernel.get(), commandList->getPrintfKernelContainer()[0]);
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, commandList->getPrintfKernelContainer().size());
}
@@ -268,13 +268,13 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithPrintfWhenAppendedToSynch
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, kernel.printPrintfOutputCalledTimes);
EXPECT_FALSE(kernel.hangDetectedPassedToPrintfOutput);
EXPECT_EQ(0u, commandList->getPrintfKernelContainer().size());
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(2u, kernel.printPrintfOutputCalledTimes);
EXPECT_FALSE(kernel.hangDetectedPassedToPrintfOutput);
@@ -296,13 +296,13 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithPrintfWhenAppendedToAsync
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(1u, kernel.printPrintfOutputCalledTimes);
EXPECT_FALSE(kernel.hangDetectedPassedToPrintfOutput);
EXPECT_EQ(0u, commandList->getPrintfKernelContainer().size());
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(2u, kernel.printPrintfOutputCalledTimes);
EXPECT_FALSE(kernel.hangDetectedPassedToPrintfOutput);
@@ -329,13 +329,13 @@ HWTEST_F(CommandListAppendLaunchKernel, givenKernelWithPrintfWhenAppendToSynchro
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_DEVICE_LOST, result);
EXPECT_EQ(1u, kernel.printPrintfOutputCalledTimes);
EXPECT_TRUE(kernel.hangDetectedPassedToPrintfOutput);
EXPECT_EQ(0u, commandList->getPrintfKernelContainer().size());
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_DEVICE_LOST, result);
EXPECT_EQ(2u, kernel.printPrintfOutputCalledTimes);
EXPECT_TRUE(kernel.hangDetectedPassedToPrintfOutput);
@@ -360,7 +360,7 @@ HWTEST_F(CommandListAppendLaunchKernel, WhenAppendingMultipleTimesThenSshIsNotDe
const_cast<KernelDescriptor::AddressingMode &>(kernel->getKernelDescriptor().kernelAttributes.bufferAddressingMode) = KernelDescriptor::BindfulAndStateless;
CmdListKernelLaunchParams launchParams = {};
for (size_t i = 0; i < sshHeapSize / kernelSshSize + 1; i++) {
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
}
@@ -395,7 +395,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenTimestampEventsWhenAppendingKernel
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(
kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();
@@ -490,7 +490,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelLaunchWithTSEventAndScopeFla
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(
kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();
@@ -519,7 +519,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenForcePipeControlPriorToWalkerKeyTh
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandListBase->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListBase->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandListBase->getCmdContainer().getCommandStream()->getUsed();
@@ -540,7 +540,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenForcePipeControlPriorToWalkerKeyTh
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
usedSpaceBefore = commandListWithDebugKey->getCmdContainer().getCommandStream()->getUsed();
result = commandListWithDebugKey->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListWithDebugKey->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
usedSpaceAfter = commandListWithDebugKey->getCmdContainer().getCommandStream()->getUsed();
@@ -576,7 +576,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenForcePipeControlPriorToWalkerKeyAn
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto secondBatchBufferAllocation = commandList->getCmdContainer().getCommandStream()->getGraphicsAllocation();
@@ -607,7 +607,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenCommandListWhenAppendLaunchKernelS
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
for (uint32_t i = 0; i < NEO::TimestampPacketConstants::preferredPacketCount + 4; i++) {
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
}
EXPECT_EQ(1u, event->getPacketsInUse());
@@ -632,13 +632,13 @@ HWTEST_F(CommandListAppendLaunchKernel, givenIndirectDispatchWhenAppendingThenWo
ASSERT_EQ(result, ZE_RESULT_SUCCESS);
result = commandList->appendLaunchKernelIndirect(kernel.toHandle(),
static_cast<ze_group_count_t *>(alloc),
*static_cast<ze_group_count_t *>(alloc),
nullptr, 0, nullptr, false);
EXPECT_EQ(result, ZE_RESULT_SUCCESS);
kernel.groupSize[2] = 2;
result = commandList->appendLaunchKernelIndirect(kernel.toHandle(),
static_cast<ze_group_count_t *>(alloc),
*static_cast<ze_group_count_t *>(alloc),
nullptr, 0, nullptr, false);
EXPECT_EQ(result, ZE_RESULT_SUCCESS);
@@ -730,7 +730,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenCommandListWhenResetCalledThenState
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(
kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -788,7 +788,7 @@ HWTEST_F(CommandListAppendLaunchKernel, WhenAddingKernelsThenResidencyContainerD
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
for (int i = 0; i < 4; ++i) {
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
}
@@ -826,7 +826,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenSingleValidWaitEventsThenAddSemapho
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 1, &hEventHandle, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 1, &hEventHandle, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();
@@ -884,7 +884,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenMultipleValidWaitEventsThenAddSemap
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 2, waitEvents, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 2, waitEvents, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();
@@ -905,7 +905,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenInvalidEventListWhenAppendLaunchCoo
ze_group_count_t groupCount{1, 1, 1};
ze_result_t returnValue;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 1, nullptr, false);
returnValue = commandList->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 1, nullptr, false);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, returnValue);
}
@@ -920,7 +920,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenImmediateCommandListWhenAppendLaun
cmdList.commandContainer.setImmediateCmdListCsr(device->getNEODevice()->getDefaultEngine().commandStreamReceiver);
ze_group_count_t groupCount{1, 1, 1};
ze_result_t returnValue;
returnValue = cmdList.appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
returnValue = cmdList.appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(0u, cmdList.executeCommandListImmediateCalledCount);
EXPECT_EQ(1u, cmdList.executeCommandListImmediateWithFlushTaskCalledCount);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
@@ -936,7 +936,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenImmediateCommandListWhenAppendLaun
cmdList.commandContainer.setImmediateCmdListCsr(device->getNEODevice()->getDefaultEngine().commandStreamReceiver);
ze_group_count_t groupCount{1, 1, 1};
ze_result_t returnValue;
returnValue = cmdList.appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, false);
returnValue = cmdList.appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(1u, cmdList.executeCommandListImmediateCalledCount);
EXPECT_EQ(0u, cmdList.executeCommandListImmediateWithFlushTaskCalledCount);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
@@ -964,13 +964,13 @@ HWTEST2_F(CommandListAppendLaunchKernel, whenUpdateStreamPropertiesIsCalledThenC
EXPECT_EQ(-1, commandList->finalStreamState.stateComputeMode.threadArbitrationPolicy.value);
const ze_group_count_t launchKernelArgs = {};
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(expectedThreadArbitrationPolicy, commandList->finalStreamState.stateComputeMode.threadArbitrationPolicy.value);
for (auto threadArbitrationPolicy : threadArbitrationPolicyValues) {
DebugManager.flags.OverrideThreadArbitrationPolicy.set(threadArbitrationPolicy);
commandList->reset();
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(threadArbitrationPolicy, commandList->finalStreamState.stateComputeMode.threadArbitrationPolicy.value);
}
}

44
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_2.cpp
View File

@@ -85,7 +85,7 @@ HWTEST_F(CommandListDualStorage, givenIndirectDispatchWithSharedDualStorageMemor
pThreadGroupDimensions->groupCountZ = 5;
result = commandList->appendLaunchKernelIndirect(kernel.toHandle(),
pThreadGroupDimensions,
*pThreadGroupDimensions,
nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -257,7 +257,7 @@ HWCMDTEST_F(IGFX_XE_HP_CORE, CommandListDualStorage, givenIndirectDispatchWithSh
pThreadGroupDimensions->groupCountZ = 5;
result = commandList->appendLaunchKernelIndirect(kernel.toHandle(),
pThreadGroupDimensions,
*pThreadGroupDimensions,
nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -402,7 +402,7 @@ HWTEST_F(CommandListAppendLaunchKernelSWTags, givenEnableSWTagsWhenAppendLaunchK
auto usedSpaceBefore = cmdStream->getUsed();
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = cmdStream->getUsed();
@@ -989,7 +989,7 @@ struct CmdlistAppendLaunchKernelWithImplicitArgsTests : CmdlistAppendLaunchKerne
ze_group_count_t groupCount{expectedImplicitArgs.groupCountX, expectedImplicitArgs.groupCountY, expectedImplicitArgs.groupCountZ};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
const auto &gfxCoreHelper = device->getGfxCoreHelper();
@@ -1137,7 +1137,7 @@ HWTEST_F(CmdlistAppendLaunchKernelTests, givenKernelWithoutImplicitArgsWhenAppen
ze_group_count_t groupCount = {3, 2, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto indirectHeap = commandList->getCmdContainer().getIndirectHeap(NEO::HeapType::INDIRECT_OBJECT);
@@ -1175,7 +1175,7 @@ HWTEST2_F(CmdlistAppendLaunchKernelTests, givenKernelWithScratchAndPrivateWhenAp
ze_group_count_t groupCount = {3, 2, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(commandList->getCommandListPerThreadPrivateScratchSize(), static_cast<uint32_t>(0x100));
@@ -1217,7 +1217,7 @@ HWTEST2_F(CmdlistAppendLaunchKernelTests, givenGlobalBindlessAllocatorAndKernelW
ze_group_count_t groupCount = {3, 2, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_NE(nullptr, commandList->getCmdContainer().getIndirectHeap(HeapType::SURFACE_STATE));
@@ -1258,7 +1258,7 @@ HWTEST2_F(CmdlistAppendLaunchKernelTests, givenGlobalBindlessAllocatorAndKernelW
ze_group_count_t groupCount = {3, 2, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_NE(nullptr, commandList->getCmdContainer().getIndirectHeap(HeapType::SURFACE_STATE));
@@ -1317,19 +1317,19 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenCooperativeAndNonCooperativeKernel
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = false;
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
launchParams.isCooperative = false;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
}
@@ -1350,7 +1350,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelWithSlmSizeExceedingLocalMem
::testing::internal::CaptureStderr();
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
std::string output = testing::internal::GetCapturedStderr();
@@ -1361,7 +1361,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelWithSlmSizeExceedingLocalMem
::testing::internal::CaptureStderr();
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
const char *pStr = nullptr;
std::string emptyString = "";
zeDriverGetLastErrorDescription(device->getDriverHandle(), &pStr);
@@ -1463,18 +1463,18 @@ HWTEST2_F(CommandListAppendLaunchKernel, GivenDebugToggleSetWhenUpdateStreamProp
// initial kernel with no policy preference
const ze_group_count_t launchKernelArgs = {};
pCommandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(defaultThreadArbitrationPolicy, pCommandList->finalStreamState.stateComputeMode.threadArbitrationPolicy.value);
// policy changed to non-default state
pCommandList->finalStreamState.stateComputeMode.threadArbitrationPolicy.value = nonDefaultThreadArbitrationPolicy;
// another kernel with no policy preference - do not update policy
pCommandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(nonDefaultThreadArbitrationPolicy, pCommandList->finalStreamState.stateComputeMode.threadArbitrationPolicy.value);
// another kernel with no policy preference, this time with debug toggle set - update policy back to default value
DebugManager.flags.ForceDefaultThreadArbitrationPolicyIfNotSpecified.set(true);
pCommandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(defaultThreadArbitrationPolicy, pCommandList->finalStreamState.stateComputeMode.threadArbitrationPolicy.value);
}
@@ -1502,7 +1502,7 @@ HWCMDTEST_F(IGFX_XE_HP_CORE, MultiTileCommandListAppendLaunchKernelXeHpCoreTest,
ze_group_count_t groupCount{256, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, hEventHandle, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, hEventHandle, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(4u, event->getPacketsInUse());
@@ -1536,7 +1536,7 @@ HWCMDTEST_F(IGFX_XE_HP_CORE, MultiTileCommandListAppendLaunchKernelXeHpCoreTest,
ze_group_count_t groupCount{256, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, hEventHandle, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, hEventHandle, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(4u, event->getPacketsInUse());
EXPECT_EQ(4u, commandList->partitionCount);
@@ -1563,7 +1563,7 @@ HWTEST2_F(MultiTileCommandListAppendLaunchKernelXeHpCoreTest, givenCooperativeKe
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto sizeBefore = commandListWithNonCooperativeKernel->getCmdContainer().getCommandStream()->getUsed();
CmdListKernelLaunchParams launchParams = {};
result = commandListWithNonCooperativeKernel->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = commandListWithNonCooperativeKernel->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto sizeAfter = commandListWithNonCooperativeKernel->getCmdContainer().getCommandStream()->getUsed();
GenCmdList cmdList;
@@ -1579,7 +1579,7 @@ HWTEST2_F(MultiTileCommandListAppendLaunchKernelXeHpCoreTest, givenCooperativeKe
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
sizeBefore = commandListWithCooperativeKernel->getCmdContainer().getCommandStream()->getUsed();
launchParams.isCooperative = true;
result = commandListWithCooperativeKernel->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = commandListWithCooperativeKernel->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
sizeAfter = commandListWithCooperativeKernel->getCmdContainer().getCommandStream()->getUsed();
cmdList.clear();
@@ -1604,7 +1604,7 @@ HWTEST2_F(MultiTileCommandListAppendLaunchKernelXeHpCoreTest,
auto sizeBefore = cmdStream->getUsed();
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.get(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.get(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto sizeAfter = cmdStream->getUsed();

212
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp
View File

@@ -62,7 +62,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, CommandListAppendLaunchKernel, givenFunctionWhenBind
ze_result_t returnValue;
std::unique_ptr<L0::CommandList> commandList(CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto commandStream = commandList->getCmdContainer().getCommandStream();
@@ -110,7 +110,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, CommandListAppendLaunchKernel, givenEventsWhenAppend
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(
kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();
@@ -155,10 +155,11 @@ HWCMDTEST_F(IGFX_GEN8_CORE, CommandListAppendLaunchKernel, givenAppendLaunchMult
auto commandList = std::unique_ptr<L0::CommandList>(L0::CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
const ze_kernel_handle_t launchKernels = kernel->toHandle();
uint32_t *numLaunchArgs;
const ze_group_count_t launchKernelArgs = {1, 1, 1};
ze_device_mem_alloc_desc_t deviceDesc = {};
auto result = context->allocDeviceMem(
device->toHandle(), &deviceDesc, 16384u, 4096u, reinterpret_cast<void **>(&numLaunchArgs));
result = commandList->appendLaunchMultipleKernelsIndirect(1, &launchKernels, numLaunchArgs, nullptr, nullptr, 0, nullptr, false);
result = commandList->appendLaunchMultipleKernelsIndirect(1, &launchKernels, numLaunchArgs, launchKernelArgs, nullptr, 0, nullptr, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
*numLaunchArgs = 0;
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();
@@ -184,10 +185,11 @@ HWCMDTEST_F(IGFX_GEN8_CORE, CommandListAppendLaunchKernel, givenAppendLaunchMult
const ze_kernel_handle_t launchKernels[3] = {kernel->toHandle(), kernel->toHandle(), kernel->toHandle()};
uint32_t *numLaunchArgs;
const uint32_t numKernels = 3;
const ze_group_count_t launchKernelArgs = {1, 1, 1};
ze_device_mem_alloc_desc_t deviceDesc = {};
auto result = context->allocDeviceMem(
device->toHandle(), &deviceDesc, 16384u, 4096u, reinterpret_cast<void **>(&numLaunchArgs));
result = commandList->appendLaunchMultipleKernelsIndirect(numKernels, launchKernels, numLaunchArgs, nullptr, nullptr, 0, nullptr, false);
result = commandList->appendLaunchMultipleKernelsIndirect(numKernels, launchKernels, numLaunchArgs, launchKernelArgs, nullptr, 0, nullptr, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
*numLaunchArgs = 2;
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();
@@ -235,7 +237,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenImmediateCommandListWhenAppendingL
CmdListKernelLaunchParams launchParams = {};
result = commandList0->appendLaunchKernel(
kernel->toHandle(),
&groupCount, nullptr, 0, nullptr, launchParams, false);
groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
}
@@ -263,7 +265,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenImmediateCommandListWhenAppendingL
CmdListKernelLaunchParams launchParams = {};
result = commandList0->appendLaunchKernel(
kernel->toHandle(),
&groupCount, nullptr, 1, nullptr, launchParams, false);
groupCount, nullptr, 1, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
}
@@ -289,7 +291,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenNonemptyAllocPrintfBufferKernelWhe
auto event = std::unique_ptr<Event>(Event::create<typename FamilyType::TimestampPacketType>(eventPool.get(), &eventDesc, device));
ze_group_count_t groupCount{1, 1, 1};
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, false);
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
ASSERT_NE(nullptr, event->getKernelForPrintf());
@@ -316,7 +318,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenEmptyAllocPrintfBufferKernelWhenAp
auto event = std::unique_ptr<Event>(Event::create<typename FamilyType::TimestampPacketType>(eventPool.get(), &eventDesc, device));
ze_group_count_t groupCount{1, 1, 1};
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, false);
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
ASSERT_EQ(nullptr, event->getKernelForPrintf());
@@ -349,7 +351,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenNonemptyAllocPrintfBufferKernelWhe
auto pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, event.get(), launchParams);
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, event.get(), launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
ASSERT_NE(nullptr, event->getKernelForPrintf());
@@ -381,7 +383,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenEmptyAllocPrintfBufferKernelWhenAp
auto pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, event.get(), launchParams);
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, event.get(), launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
ASSERT_EQ(nullptr, event->getKernelForPrintf());
@@ -409,7 +411,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenImmediateCommandListWhenAppendingL
result = commandList0->appendLaunchKernelIndirect(
kernel->toHandle(),
&groupCount, nullptr, 0, nullptr, false);
groupCount, nullptr, 0, nullptr, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
}
@@ -436,7 +438,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenImmediateCommandListWhenAppendingL
result = commandList0->appendLaunchKernelIndirect(
kernel->toHandle(),
&groupCount, nullptr, 1, nullptr, false);
groupCount, nullptr, 1, nullptr, false);
ASSERT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
}
@@ -460,14 +462,14 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelUsingSyncBufferWhenAppendLau
engineGroupType = gfxCoreHelper.getEngineGroupType(aub_stream::EngineType::ENGINE_CCS, EngineUsage::Cooperative, *defaultHwInfo);
}
pCommandList->initialize(device, engineGroupType, 0u);
auto result = pCommandList->appendLaunchCooperativeKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, false);
auto result = pCommandList->appendLaunchCooperativeKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, engineGroupType, 0u);
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
{
@@ -475,7 +477,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelUsingSyncBufferWhenAppendLau
usesSyncBuffer = {};
pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
}
{
@@ -485,13 +487,13 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelUsingSyncBufferWhenAppendLau
groupCountX = maximalNumberOfWorkgroupsAllowed + 1;
pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, engineGroupType, 0u);
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
}
{
VariableBackup<bool> cooperative{&launchParams.isCooperative};
cooperative = false;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
}
}
@@ -511,12 +513,12 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenDisableOverdispatchPropertyWhenUpd
EXPECT_EQ(expectedDisableOverdispatch, pCommandList->requiredStreamState.frontEndState.disableOverdispatch.value);
EXPECT_EQ(expectedDisableOverdispatch, pCommandList->finalStreamState.frontEndState.disableOverdispatch.value);
const ze_group_count_t launchKernelArgs = {};
pCommandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
const ze_group_count_t launchKernelArgs = {1, 1, 1};
pCommandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(expectedDisableOverdispatch, pCommandList->requiredStreamState.frontEndState.disableOverdispatch.value);
EXPECT_EQ(expectedDisableOverdispatch, pCommandList->finalStreamState.frontEndState.disableOverdispatch.value);
pCommandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_EQ(expectedDisableOverdispatch, pCommandList->requiredStreamState.frontEndState.disableOverdispatch.value);
EXPECT_EQ(expectedDisableOverdispatch, pCommandList->finalStreamState.frontEndState.disableOverdispatch.value);
}
@@ -530,7 +532,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenCooperativeKernelWhenAppendLaunchC
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = false;
auto result = pCommandList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
auto result = pCommandList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_TRUE(pCommandList->containsAnyKernel);
EXPECT_FALSE(pCommandList->containsCooperativeKernelsFlag);
@@ -538,7 +540,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenCooperativeKernelWhenAppendLaunchC
pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_TRUE(pCommandList->containsAnyKernel);
EXPECT_TRUE(pCommandList->containsCooperativeKernelsFlag);
@@ -555,19 +557,19 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenAnyCooperativeKernelAndMixingAllow
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = false;
auto result = pCommandList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
auto result = pCommandList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_TRUE(pCommandList->containsAnyKernel);
EXPECT_FALSE(pCommandList->containsCooperativeKernelsFlag);
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_TRUE(pCommandList->containsAnyKernel);
EXPECT_TRUE(pCommandList->containsCooperativeKernelsFlag);
launchParams.isCooperative = false;
result = pCommandList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_TRUE(pCommandList->containsAnyKernel);
EXPECT_TRUE(pCommandList->containsCooperativeKernelsFlag);
@@ -587,19 +589,19 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenCooperativeAndNonCooperativeKernel
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = false;
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
auto result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
pCommandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
pCommandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
launchParams.isCooperative = false;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
}
@@ -657,7 +659,7 @@ HWTEST_F(CommandListAppendLaunchKernel, givenInvalidKernelWhenAppendingThenRetur
ze_group_count_t groupCount{8, 1, 1};
CmdListKernelLaunchParams launchParams = {};
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, returnValue);
}
@@ -844,7 +846,7 @@ HWTEST2_F(InOrderCmdListTests, givenNotSignaledInOrderEventWhenAddedToWaitListTh
auto handle = event->toHandle();
returnValue = immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, &handle, launchParams, false);
returnValue = immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, &handle, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, returnValue);
}
@@ -862,9 +864,9 @@ HWTEST2_F(InOrderCmdListTests, givenCmdListsWhenDispatchingThenUseInternalTaskCo
auto ultCsr = static_cast<UltCommandStreamReceiver<FamilyType> *>(device->getNEODevice()->getDefaultEngine().commandStreamReceiver);
immCmdList0->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList0->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList1->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList1->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(1u, immCmdList0->cmdQImmediate->getTaskCount());
EXPECT_EQ(2u, immCmdList1->cmdQImmediate->getTaskCount());
@@ -907,7 +909,7 @@ HWTEST2_F(InOrderCmdListTests, givenDebugFlagSetWhenEventHostSyncCalledThenCallW
auto eventPool = createEvents<FamilyType>(2, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
events[0]->inOrderAllocationOffset = 123;
@@ -950,7 +952,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenResetEventCalledThenResetEven
auto eventPool = createEvents<FamilyType>(3, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(MemoryConstants::pageSize64k, immCmdList->inOrderDependencyCounterAllocation->getUnderlyingBufferSize());
@@ -973,7 +975,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWheUsingRegularEventThenDontSetIn
auto eventPool = createEvents<FamilyType>(1, false);
events[0]->inOrderExecEvent = false;
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
EXPECT_FALSE(events[0]->inOrderExecEvent);
EXPECT_EQ(events[0]->inOrderExecSignalValue, 0u);
@@ -991,11 +993,11 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenSubmittingThenProgramSemaphor
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(
@@ -1028,7 +1030,7 @@ HWTEST2_F(InOrderCmdListTests, givenDebugFlagSetWhenDispatchingSemaphoreThenProg
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -1061,7 +1063,7 @@ HWTEST2_F(InOrderCmdListTests, givenTimestmapEventWhenProgrammingBarrierThenDont
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -1092,7 +1094,7 @@ HWTEST2_F(InOrderCmdListTests, givenDebugFlagSetWhenDispatchingStoreDataImmThenP
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -1136,7 +1138,7 @@ HWTEST2_F(InOrderCmdListTests, givenDebugFlagSetAsMaskWhenDispatchingStoreDataIm
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -1181,11 +1183,11 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenWaitingForEventFromPreviousAp
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, &eventHandle, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, &eventHandle, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(
@@ -1208,10 +1210,10 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenWaitingForEventFromAfterReset
auto eventPool = createEvents<FamilyType>(1, false);
auto eventHandle = events[0]->toHandle();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
events[0]->reset();
auto retValue = immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, &eventHandle, launchParams, false);
auto retValue = immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, &eventHandle, launchParams, false);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, retValue);
}
@@ -1236,16 +1238,16 @@ HWTEST2_F(InOrderCmdListTests, givenMultipleAllocationOwnerWhenUsingEventsThenSe
validateNumOwners(1);
immCmdList0->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle0, 0, nullptr, launchParams, false);
immCmdList0->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle0, 0, nullptr, launchParams, false);
validateNumOwners(2);
immCmdList0->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle0, 0, nullptr, launchParams, false);
immCmdList0->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle0, 0, nullptr, launchParams, false);
validateNumOwners(2);
immCmdList0->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle1, 0, nullptr, launchParams, false);
immCmdList0->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle1, 0, nullptr, launchParams, false);
validateNumOwners(3);
immCmdList1->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle0, 0, nullptr, launchParams, false);
immCmdList1->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle0, 0, nullptr, launchParams, false);
validateNumOwners(2);
events[1]->reset();
@@ -1271,17 +1273,17 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderEventModeWhenSubmittingThenProgramSem
auto event0Handle = events[0]->toHandle();
auto event1Handle = events[1]->toHandle();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, event0Handle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, event0Handle, 0, nullptr, launchParams, false);
immCmdList2->appendLaunchKernel(kernel->toHandle(), &groupCount, event1Handle, 0, nullptr, launchParams, false);
immCmdList2->appendLaunchKernel(kernel->toHandle(), groupCount, event1Handle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
ze_event_handle_t waitlist[] = {event0Handle, event1Handle};
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 2, waitlist, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 2, waitlist, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(
@@ -1318,13 +1320,13 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderEventModeWhenWaitingForEventFromPrevi
auto event0Handle = events[0]->toHandle();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, event0Handle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, event0Handle, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, &event0Handle, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, &event0Handle, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(
@@ -1356,11 +1358,11 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderEventModeWhenSubmittingFromDifferentC
auto ultCsr = static_cast<UltCommandStreamReceiver<FamilyType> *>(device->getNEODevice()->getDefaultEngine().commandStreamReceiver);
ultCsr->storeMakeResidentAllocations = true;
immCmdList1->appendLaunchKernel(kernel->toHandle(), &groupCount, event0Handle, 0, nullptr, launchParams, false);
immCmdList1->appendLaunchKernel(kernel->toHandle(), groupCount, event0Handle, 0, nullptr, launchParams, false);
EXPECT_EQ(1u, ultCsr->makeResidentAllocations[immCmdList1->inOrderDependencyCounterAllocation]);
immCmdList2->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, &event0Handle, launchParams, false);
immCmdList2->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, &event0Handle, launchParams, false);
EXPECT_EQ(2u, ultCsr->makeResidentAllocations[immCmdList1->inOrderDependencyCounterAllocation]);
@@ -1389,7 +1391,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderEventModeWhenSubmittingThenClearEvent
events[0]->csrs.clear();
events[0]->csrs.push_back(&tempCsr);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(1u, events[0]->csrs.size());
EXPECT_EQ(device->getNEODevice()->getDefaultEngine().commandStreamReceiver, events[0]->csrs[0]);
@@ -1406,11 +1408,11 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenDispatchingThenHandleDependen
auto ultCsr = static_cast<UltCommandStreamReceiver<FamilyType> *>(device->getNEODevice()->getDefaultEngine().commandStreamReceiver);
ultCsr->storeMakeResidentAllocations = true;
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(1u, immCmdList->inOrderDependencyCounter);
EXPECT_EQ(1u, ultCsr->makeResidentAllocations[immCmdList->inOrderDependencyCounterAllocation]);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(2u, immCmdList->inOrderDependencyCounter);
EXPECT_EQ(2u, ultCsr->makeResidentAllocations[immCmdList->inOrderDependencyCounterAllocation]);
}
@@ -1422,7 +1424,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenAddingRelaxedOrderingEventsTh
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -1459,7 +1461,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingWalkerThenSignalSy
auto eventPool = createEvents<FamilyType>(1, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
{
@@ -1479,7 +1481,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingWalkerThenSignalSy
auto offset = cmdStream->getUsed();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
{
@@ -1739,7 +1741,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingRegularEventThenCl
events[0]->signalScope = 0;
events[0]->inOrderExecEvent = false;
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList, cmdStream->getCpuBase(), cmdStream->getUsed()));
@@ -1791,7 +1793,7 @@ HWTEST2_F(InOrderCmdListTests, givenHostVisibleEventOnLatestFlushWhenCallingSync
EXPECT_FALSE(immCmdList->latestFlushIsHostVisible);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
EXPECT_FALSE(immCmdList->latestFlushIsHostVisible);
EXPECT_EQ(0u, immCmdList->synchronizeInOrderExecutionCalled);
@@ -1802,7 +1804,7 @@ HWTEST2_F(InOrderCmdListTests, givenHostVisibleEventOnLatestFlushWhenCallingSync
EXPECT_EQ(1u, ultCsr->waitForCompletionWithTimeoutTaskCountCalled);
events[0]->signalScope = ZE_EVENT_SCOPE_FLAG_HOST;
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
EXPECT_TRUE(immCmdList->latestFlushIsHostVisible);
immCmdList->hostSynchronize(0, 1, false);
@@ -1851,7 +1853,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingWalkerThenProgramP
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList, cmdStream->getCpuBase(), cmdStream->getUsed()));
@@ -1939,7 +1941,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingAppendSignalEventT
auto eventPool = createEvents<FamilyType>(1, true);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -1997,7 +1999,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingNonKernelAppendThe
uint8_t ptr[64] = {};
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
uint32_t inOrderCounter = 1;
@@ -2087,7 +2089,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderRegularCmdListWhenProgrammingNonKerne
uint64_t inOrderSyncVa = regularCmdList->inOrderDependencyCounterAllocation->getGpuAddress();
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto verifySemaphore = [&inOrderSyncVa](const GenCmdList::iterator &iterator, uint64_t waitValue) {
auto semaphoreCmd = genCmdCast<MI_SEMAPHORE_WAIT *>(*iterator);
@@ -2206,9 +2208,9 @@ HWTEST2_F(InOrderCmdListTests, givenImmediateEventWhenWaitingFromRegularCmdListT
auto eventPool = createEvents<FamilyType>(1, false);
auto eventHandle = events[0]->toHandle();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, &eventHandle, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, &eventHandle, launchParams, false);
ASSERT_EQ(1u, regularCmdList->inOrderPatchCmds.size());
@@ -2547,8 +2549,8 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingAppendWaitOnEvents
auto eventHandle = events[0]->toHandle();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -2591,8 +2593,8 @@ HWTEST2_F(InOrderCmdListTests, givenRegularInOrderCmdListWhenProgrammingAppendWa
auto eventHandle = events[0]->toHandle();
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -2637,7 +2639,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingCounterWithOverflo
uint64_t baseGpuVa = immCmdList->inOrderDependencyCounterAllocation->getGpuAddress();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList, cmdStream->getCpuBase(), cmdStream->getUsed()));
@@ -2741,7 +2743,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingAppendBarrierWithW
auto eventHandle = events[0]->toHandle();
immCmdList1->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList1->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -2771,7 +2773,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingAppendBarrierWitho
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(1u, immCmdList->inOrderDependencyCounter);
@@ -2799,9 +2801,9 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingAppendBarrierWithD
auto eventPool = createEvents<FamilyType>(3, false);
immCmdList1->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList2->appendLaunchKernel(kernel->toHandle(), &groupCount, events[1]->toHandle(), 0, nullptr, launchParams, false);
immCmdList2->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList1->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList2->appendLaunchKernel(kernel->toHandle(), groupCount, events[1]->toHandle(), 0, nullptr, launchParams, false);
immCmdList2->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
auto offset = cmdStream->getUsed();
@@ -2835,7 +2837,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingAppendBarrierWitho
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(1u, immCmdList->inOrderDependencyCounter);
@@ -2871,7 +2873,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenProgrammingAppendBarrierWitho
auto cmdStream = immCmdList->getCmdContainer().getCommandStream();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(1u, immCmdList->inOrderDependencyCounter);
@@ -2910,7 +2912,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenCallingSyncThenHandleCompleti
auto eventPool = createEvents<FamilyType>(1, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
auto hostAddress = static_cast<uint64_t *>(ptrOffset(immCmdList->inOrderDependencyCounterAllocation->getUnderlyingBuffer(), counterOffset));
*hostAddress = 0;
@@ -2967,7 +2969,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenCallingSyncThenHandleCompleti
EXPECT_EQ(1u, *hostAddress);
}
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
*ultCsr->getTagAddress() = ultCsr->taskCount - 1;
EXPECT_EQ(ZE_RESULT_NOT_READY, immCmdList->hostSynchronize(0, ultCsr->taskCount, true));
@@ -2999,7 +3001,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenDoingCpuCopyThenSynchronize,
}
};
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
events[0]->setIsCompleted();
ultCsr->waitForCompletionWithTimeoutTaskCountCalled = 0;
@@ -3044,7 +3046,7 @@ HWTEST2_F(InOrderCmdListTests, givenInOrderModeWhenGpuHangDetectedInCpuCopyPathT
auto hostAddress = static_cast<uint64_t *>(immCmdList->inOrderDependencyCounterAllocation->getUnderlyingBuffer());
*hostAddress = 0;
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
void *deviceAlloc = nullptr;
ze_device_mem_alloc_desc_t deviceDesc = {};
@@ -3182,11 +3184,11 @@ HWTEST2_F(MultiTileInOrderCmdListTests, givenMultiTileInOrderModeWhenProgramming
auto eventPool = createEvents<FamilyType>(1, false);
auto eventHandle = events[0]->toHandle();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
size_t offset = cmdStream->getUsed();
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 1, &eventHandle, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 1, &eventHandle, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList,
@@ -3236,7 +3238,7 @@ HWTEST2_F(MultiTileInOrderCmdListTests, givenMultiTileInOrderModeWhenCallingSync
auto eventPool = createEvents<FamilyType>(1, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, events[0]->toHandle(), 0, nullptr, launchParams, false);
auto hostAddress0 = static_cast<uint64_t *>(immCmdList->inOrderDependencyCounterAllocation->getUnderlyingBuffer());
auto hostAddress1 = ptrOffset(hostAddress0, sizeof(uint64_t));
@@ -3278,7 +3280,7 @@ HWTEST2_F(MultiTileInOrderCmdListTests, givenMultiTileInOrderModeWhenProgramming
auto eventHandle = events[0]->toHandle();
events[0]->signalScope = 0;
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList,
@@ -3330,7 +3332,7 @@ HWTEST2_F(MultiTileInOrderCmdListTests, givenMultiTileInOrderModeWhenProgramming
immCmdList->signalAllEventPackets = true;
events[0]->maxPacketCount = 4;
immCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, eventHandle, 0, nullptr, launchParams, false);
immCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, eventHandle, 0, nullptr, launchParams, false);
GenCmdList cmdList;
ASSERT_TRUE(FamilyType::PARSE::parseCommandBuffer(cmdList,
@@ -3392,8 +3394,8 @@ HWTEST2_F(MultiTileInOrderCmdListTests, whenUsingRegularCmdListThenAddWalkerToPa
size_t offset = cmdStream->getUsed();
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(4u, regularCmdList->inOrderPatchCmds.size()); // Walker + 2x Semaphore + Walker
@@ -3704,8 +3706,8 @@ HWTEST2_F(InOrderRegularCmdListTests, whenUsingRegularCmdListThenAddWalkerToPatc
size_t offset = cmdStream->getUsed();
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(3u, regularCmdList->inOrderPatchCmds.size()); // Walker + Semaphore + Walker
@@ -3772,7 +3774,7 @@ HWTEST2_F(InOrderRegularCmdListTests, givenInOrderModeWhenDispatchingRegularCmdL
size_t offset = cmdStream->getUsed();
EXPECT_EQ(0u, regularCmdList->inOrderDependencyCounter);
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(1u, regularCmdList->inOrderDependencyCounter);
{
@@ -3797,7 +3799,7 @@ HWTEST2_F(InOrderRegularCmdListTests, givenInOrderModeWhenDispatchingRegularCmdL
offset = cmdStream->getUsed();
regularCmdList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
regularCmdList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(2u, regularCmdList->inOrderDependencyCounter);
{
@@ -3999,7 +4001,7 @@ HWTEST_F(CommandListAppendLaunchKernelWithImplicitArgs, givenIndirectDispatchWit
ASSERT_EQ(result, ZE_RESULT_SUCCESS);
result = commandList->appendLaunchKernelIndirect(kernel.toHandle(),
static_cast<ze_group_count_t *>(alloc),
*static_cast<ze_group_count_t *>(alloc),
nullptr, 0, nullptr, false);
EXPECT_EQ(result, ZE_RESULT_SUCCESS);
auto heap = commandList->getCmdContainer().getIndirectHeap(HeapType::INDIRECT_OBJECT);
@@ -4173,13 +4175,13 @@ HWTEST2_F(MultiTileImmediateCommandListAppendLaunchKernelXeHpCoreTest, givenImpl
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
CmdListKernelLaunchParams launchParams = {};
result = immediateCmdList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = immediateCmdList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto cmdStream = immediateCmdList->getCmdContainer().getCommandStream();
auto sizeBefore = cmdStream->getUsed();
result = immediateCmdList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = immediateCmdList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto sizeAfter = cmdStream->getUsed();
@@ -4238,7 +4240,7 @@ HWTEST2_F(MultiTileImmediateCommandListAppendLaunchKernelXeHpCoreTest, givenImpl
auto sizeBefore = cmdStream->getUsed();
CmdListKernelLaunchParams launchParams = {};
result = immediateCmdList->appendLaunchKernelWithParams(kernel.get(), &groupCount, nullptr, launchParams);
result = immediateCmdList->appendLaunchKernelWithParams(kernel.get(), groupCount, nullptr, launchParams);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto sizeAfter = cmdStream->getUsed();

6
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_wait_on_events.cpp
View File

@@ -686,7 +686,7 @@ HWTEST_TEMPLATED_F(TbxImmediateCommandListTest, givenTbxModeOnFlushTaskImmediate
auto eventHandle = event->toHandle();
ze_group_count_t group = {1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandListImmediate->appendLaunchKernel(kernel->toHandle(), &group, nullptr, 1, &eventHandle, launchParams, false);
commandListImmediate->appendLaunchKernel(kernel->toHandle(), group, nullptr, 1, &eventHandle, launchParams, false);
EXPECT_EQ(1u, ultCsr.downloadAllocationsCalledCount);
}
@@ -696,7 +696,7 @@ HWTEST_TEMPLATED_F(TbxImmediateCommandListTest, givenTbxModeOnFlushTaskImmediate
auto eventHandle = event->toHandle();
ze_group_count_t group = {1, 1, 1};
commandListImmediate->appendLaunchKernelIndirect(kernel->toHandle(), &group, nullptr, 1, &eventHandle, false);
commandListImmediate->appendLaunchKernelIndirect(kernel->toHandle(), group, nullptr, 1, &eventHandle, false);
EXPECT_EQ(1u, ultCsr.downloadAllocationsCalledCount);
}
@@ -873,7 +873,7 @@ HWTEST_TEMPLATED_F(TbxImmediateCommandListTest, givenTbxModeOnFlushTaskImmediate
ze_group_count_t groupCount{1, 1, 1};
auto eventHandle = event->toHandle();
commandListImmediate->appendLaunchCooperativeKernel(kernel->toHandle(), &groupCount, nullptr, 1, &eventHandle, false);
commandListImmediate->appendLaunchCooperativeKernel(kernel->toHandle(), groupCount, nullptr, 1, &eventHandle, false);
EXPECT_EQ(1u, ultCsr.downloadAllocationsCalledCount);
}

82
level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_xehp_and_later.cpp
View File

@@ -245,23 +245,23 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenVariousKernelsWhenUpdateStreamProp
int32_t expectedDispatchAllWalkerEnable = productHelper.isComputeDispatchAllWalkerEnableInCfeStateRequired(device->getHwInfo()) ? 0 : -1;
const ze_group_count_t launchKernelArgs = {};
pCommandList->updateStreamProperties(defaultKernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, launchKernelArgs, false);
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->requiredStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->finalStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
EXPECT_EQ(0u, pCommandList->commandsToPatch.size());
pCommandList->reset();
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
expectedDispatchAllWalkerEnable = expectedDispatchAllWalkerEnable != -1 ? 1 : expectedDispatchAllWalkerEnable;
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->requiredStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->finalStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
EXPECT_EQ(0u, pCommandList->commandsToPatch.size());
pCommandList->reset();
pCommandList->updateStreamProperties(defaultKernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
expectedDispatchAllWalkerEnable = expectedDispatchAllWalkerEnable != -1 ? 0 : expectedDispatchAllWalkerEnable;
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->requiredStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
expectedDispatchAllWalkerEnable = expectedDispatchAllWalkerEnable != -1 ? 1 : expectedDispatchAllWalkerEnable;
@@ -270,18 +270,18 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenVariousKernelsWhenUpdateStreamProp
EXPECT_EQ(expectedCommandsToPatch, pCommandList->commandsToPatch.size());
pCommandList->reset();
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->requiredStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->finalStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
expectedCommandsToPatch = expectedCommandsToPatch != 0 ? 2 : 0;
EXPECT_EQ(expectedCommandsToPatch, pCommandList->commandsToPatch.size());
pCommandList->reset();
pCommandList->updateStreamProperties(defaultKernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
expectedDispatchAllWalkerEnable = expectedDispatchAllWalkerEnable != -1 ? 0 : expectedDispatchAllWalkerEnable;
EXPECT_EQ(expectedDispatchAllWalkerEnable, pCommandList->requiredStreamState.frontEndState.computeDispatchAllWalkerEnable.value);
expectedDispatchAllWalkerEnable = expectedDispatchAllWalkerEnable != -1 ? 1 : expectedDispatchAllWalkerEnable;
@@ -312,14 +312,14 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenVariousKernelsAndPatchingDisallowe
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
const ze_group_count_t launchKernelArgs = {};
pCommandList->updateStreamProperties(defaultKernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
EXPECT_EQ(0u, pCommandList->commandsToPatch.size());
pCommandList->reset();
DebugManager.flags.AllowPatchingVfeStateInCommandLists.set(1);
pCommandList->updateStreamProperties(defaultKernel, false, &launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, &launchKernelArgs, false);
pCommandList->updateStreamProperties(defaultKernel, false, launchKernelArgs, false);
pCommandList->updateStreamProperties(cooperativeKernel, true, launchKernelArgs, false);
const auto &productHelper = device->getProductHelper();
size_t expectedCmdsToPatch = productHelper.isComputeDispatchAllWalkerEnableInCfeStateRequired(device->getHwInfo()) ? 1 : 0;
@@ -395,7 +395,7 @@ struct CommandListAppendLaunchKernelCompactL3FlushEventFixture : public ModuleFi
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_EQ(arg.expectedPacketsInUse, event->getPacketsInUse());
EXPECT_EQ(arg.expectedKernelCount, event->getKernelCount());
@@ -639,7 +639,7 @@ struct CommandListSignalAllEventPacketFixture : public ModuleFixture {
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t sizeBefore = cmdStream->getUsed();
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
size_t sizeAfter = cmdStream->getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1610,11 +1610,11 @@ HWTEST2_F(CommandListAppendLaunchRayTracingKernelTest, givenKernelUsingRayTracin
neoDevice->rtMemoryBackedBuffer = nullptr;
CmdListKernelLaunchParams launchParams = {};
result = pCommandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = pCommandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_NE(ZE_RESULT_SUCCESS, result);
neoDevice->rtMemoryBackedBuffer = buffer1;
result = pCommandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = pCommandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
neoDevice->rtMemoryBackedBuffer = nullptr;
@@ -1695,7 +1695,7 @@ HWTEST2_F(RayTracingCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t regularSizeBefore = cmdStreamRegular.getUsed();
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t regularSizeAfter = cmdStreamRegular.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1745,7 +1745,7 @@ HWTEST2_F(RayTracingCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t regularSizeBefore = cmdStreamRegular.getUsed();
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t regularSizeAfter = cmdStreamRegular.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1774,7 +1774,7 @@ HWTEST2_F(RayTracingCmdListTest,
size_t immediateSizeBefore = cmdStreamImmediate.getUsed();
size_t csrBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrAfter = csrStream.getUsed();
size_t immediateSizeAfter = cmdStreamImmediate.getUsed();
@@ -1805,7 +1805,7 @@ HWTEST2_F(RayTracingCmdListTest,
size_t csrBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrAfter = csrStream.getUsed();
size_t immediateSizeAfter = cmdStreamImmediate.getUsed();
@@ -1819,7 +1819,7 @@ HWTEST2_F(RayTracingCmdListTest,
immediateSizeBefore = cmdStreamImmediate.getUsed();
csrBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
csrAfter = csrStream.getUsed();
immediateSizeAfter = cmdStreamImmediate.getUsed();
@@ -1850,7 +1850,7 @@ HWTEST2_F(RayTracingCmdListTest,
size_t csrBefore = csrStream.getUsed();
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
size_t csrAfter = csrStream.getUsed();
size_t immediateSizeAfter = cmdStreamImmediate.getUsed();
@@ -1866,7 +1866,7 @@ HWTEST2_F(RayTracingCmdListTest,
auto &cmdStreamRegular = *containerRegular.getCommandStream();
size_t regularSizeBefore = cmdStreamRegular.getUsed();
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t regularSizeAfter = cmdStreamRegular.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1903,7 +1903,7 @@ HWTEST2_F(ImmediateFlushTaskGlobalStatelessCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1932,7 +1932,7 @@ HWTEST2_F(ImmediateFlushTaskGlobalStatelessCmdListTest,
EXPECT_TRUE(csrImmediate.isMadeResident(globalSurfaceHeap->getGraphicsAllocation()));
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1956,7 +1956,7 @@ HWTEST2_F(ImmediateFlushTaskGlobalStatelessCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1982,7 +1982,7 @@ HWTEST2_F(ImmediateFlushTaskGlobalStatelessCmdListTest,
kernel->kernelRequiresUncachedMocsCount++;
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2013,7 +2013,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2054,7 +2054,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
}
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2078,7 +2078,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2112,7 +2112,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
kernel->kernelRequiresUncachedMocsCount++;
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2142,7 +2142,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2162,7 +2162,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
mockKernelImmData->kernelDescriptor->kernelAttributes.perThreadScratchSize[0] = 0x100;
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2231,7 +2231,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2263,7 +2263,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
EXPECT_EQ(this->dshRequired, sbaCmd->getDynamicStateBaseAddressModifyEnable());
EXPECT_EQ(dsShareBaseAddress, sbaCmd->getDynamicStateBaseAddress());
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2287,7 +2287,7 @@ HWTEST2_F(ImmediateFlushTaskCsrSharedHeapCmdListTest,
EXPECT_EQ(dsRegularBaseAddress, static_cast<uint64_t>(csrBaseAddressState.dynamicStateBaseAddress.value));
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2321,7 +2321,7 @@ HWTEST2_F(ImmediateFlushTaskPrivateHeapCmdListTest,
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
size_t csrUsedBefore = csrStream.getUsed();
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
size_t csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -2362,7 +2362,7 @@ HWTEST2_F(ImmediateFlushTaskPrivateHeapCmdListTest,
}
csrUsedBefore = csrStream.getUsed();
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
result = commandListImmediate->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
csrUsedAfter = csrStream.getUsed();
EXPECT_EQ(ZE_RESULT_SUCCESS, result);

6
level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_1.cpp
View File

@@ -333,7 +333,7 @@ HWTEST_F(CommandQueueCreate, given100CmdListsWhenExecutingThenCommandStreamIsNot
ze_group_count_t dispatchKernelArguments{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel.toHandle(), &dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
commandList->close();
@@ -375,7 +375,7 @@ HWTEST2_F(CommandQueueCreate, givenOutOfHostMemoryErrorFromSubmitBatchBufferWhen
ze_group_count_t dispatchFunctionArguments{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel.toHandle(), &dispatchFunctionArguments, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), dispatchFunctionArguments, nullptr, 0, nullptr, launchParams, false);
commandList->close();
ze_command_list_handle_t cmdListHandles[1] = {commandList->toHandle()};
@@ -399,7 +399,7 @@ HWTEST2_F(CommandQueueCreate, givenGpuHangInReservingLinearStreamWhenExecutingCo
ze_group_count_t dispatchKernelArguments{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel.toHandle(), &dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
commandList->close();
ze_command_list_handle_t cmdListHandles[1] = {commandList->toHandle()};

2
level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_2.cpp
View File

@@ -737,7 +737,7 @@ HWTEST2_F(DeviceWithDualStorage, givenCmdListWithAppendedKernelAndUsmTransferAnd
ze_group_count_t dispatchKernelArguments{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
commandList->appendLaunchKernel(kernel.toHandle(), &dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
commandList->appendLaunchKernel(kernel.toHandle(), dispatchKernelArguments, nullptr, 0, nullptr, launchParams, false);
auto deviceImp = static_cast<DeviceImp *>(device);
commandList->close();

40
level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_3.cpp
View File

@@ -435,7 +435,7 @@ HWTEST_F(CommandQueueIndirectAllocations, givenDebugModeToTreatIndirectAllocatio
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(),
&groupCount,
groupCount,
nullptr, 0, nullptr,
launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
@@ -498,7 +498,7 @@ HWTEST_F(CommandQueueIndirectAllocations, givenDeviceThatSupportsSubmittingIndir
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(),
&groupCount,
groupCount,
nullptr, 0, nullptr,
launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
@@ -565,7 +565,7 @@ HWTEST_F(CommandQueueIndirectAllocations, givenDeviceThatSupportsSubmittingIndir
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(),
&groupCount,
groupCount,
nullptr, 0, nullptr,
launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
@@ -617,7 +617,7 @@ HWTEST_F(CommandQueueIndirectAllocations, givenImmediateCommandListAndFlushTaskW
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(),
&groupCount,
groupCount,
nullptr, 0, nullptr,
launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
@@ -1081,57 +1081,57 @@ HWTEST2_F(CommandQueueTest, whenExecuteCommandListsIsCalledThenCorrectSizeOfFron
commandListA->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
CmdListKernelLaunchParams launchParams = {};
commandListA->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandListA->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
commandListA->close();
auto commandListBB = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandListBB->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
launchParams.isCooperative = true;
commandListBB->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListBB->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListBB->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
commandListBB->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
commandListBB->close();
auto commandListAB = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandListAB->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
launchParams.isCooperative = false;
commandListAB->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandListAB->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
launchParams.isCooperative = true;
commandListAB->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListAB->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
commandListAB->close();
auto commandListBA = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandListBA->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
commandListBA->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListBA->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
launchParams.isCooperative = false;
commandListBA->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandListBA->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
commandListBA->close();
auto commandListBAB = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandListBAB->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
launchParams.isCooperative = true;
commandListBAB->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListBAB->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
launchParams.isCooperative = false;
commandListBAB->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandListBAB->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
launchParams.isCooperative = true;
commandListBAB->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListBAB->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
commandListBAB->close();
auto commandListAAB = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandListAAB->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
launchParams.isCooperative = false;
commandListAAB->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandListAAB->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandListAAB->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
commandListAAB->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
launchParams.isCooperative = true;
commandListAAB->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListAAB->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
commandListAAB->close();
auto commandListEmpty = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
@@ -1231,7 +1231,7 @@ HWTEST2_F(CommandQueueTest, givenRegularKernelScheduledAsCooperativeWhenExecuteC
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = true;
commandList->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandList->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
commandList->close();
EXPECT_EQ(-1, csr->getStreamProperties().frontEndState.computeDispatchAllWalkerEnable.value);
@@ -1275,14 +1275,14 @@ HWTEST2_F(CommandQueueTest, givenTwoCommandQueuesUsingOneCsrWhenExecuteCommandLi
commandListA->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
CmdListKernelLaunchParams launchParams = {};
commandListA->appendLaunchKernelWithParams(&defaultKernel, &threadGroupDimensions, nullptr, launchParams);
commandListA->appendLaunchKernelWithParams(&defaultKernel, threadGroupDimensions, nullptr, launchParams);
commandListA->close();
auto commandListB = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
commandListB->initialize(device, NEO::EngineGroupType::CooperativeCompute, 0u);
launchParams.isCooperative = true;
commandListB->appendLaunchKernelWithParams(&cooperativeKernel, &threadGroupDimensions, nullptr, launchParams);
commandListB->appendLaunchKernelWithParams(&cooperativeKernel, threadGroupDimensions, nullptr, launchParams);
commandListB->close();
ze_command_list_handle_t commandListsA[] = {commandListA->toHandle()};

2
level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_debugger.cpp
View File

@@ -51,7 +51,7 @@ HWTEST_F(L0CmdQueueDebuggerTest, givenDebuggingEnabledWhenCmdListRequiringSbaPro
auto usedSpaceBefore = cmdStream.getUsed();
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
commandList->close();

4
level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist.cpp
View File

@@ -594,7 +594,7 @@ HWTEST2_F(CommandQueueExecuteCommandListsImplicitScalingDisabled, givenCommandLi
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = true;
pCommandListWithCooperativeKernels->appendLaunchKernelWithParams(&kernel, &threadGroupDimensions, nullptr, launchParams);
pCommandListWithCooperativeKernels->appendLaunchKernelWithParams(&kernel, threadGroupDimensions, nullptr, launchParams);
ze_command_list_handle_t commandListCooperative[] = {pCommandListWithCooperativeKernels->toHandle()};
auto result = pCommandQueue->executeCommandLists(1, commandListCooperative, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
@@ -604,7 +604,7 @@ HWTEST2_F(CommandQueueExecuteCommandListsImplicitScalingDisabled, givenCommandLi
pCommandListWithNonCooperativeKernels->initialize(&device, NEO::EngineGroupType::Compute, 0u);
launchParams.isCooperative = false;
pCommandListWithNonCooperativeKernels->appendLaunchKernelWithParams(&kernel, &threadGroupDimensions, nullptr, launchParams);
pCommandListWithNonCooperativeKernels->appendLaunchKernelWithParams(&kernel, threadGroupDimensions, nullptr, launchParams);
ze_command_list_handle_t commandListNonCooperative[] = {pCommandListWithNonCooperativeKernels->toHandle()};
result = pCommandQueue->executeCommandLists(1, commandListNonCooperative, nullptr, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);

6
level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist_2.cpp
View File

@@ -380,7 +380,7 @@ struct PauseOnGpuFixture : public Test<ModuleFixture> {
}
void enqueueKernel() {
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -422,7 +422,7 @@ struct PauseOnGpuTests : public PauseOnGpuFixture {
}
void enqueueKernel() {
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
result = commandList->close();
@@ -587,7 +587,7 @@ struct PauseOnGpuWithImmediateCommandListTests : public PauseOnGpuFixture {
}
void enqueueKernel() {
auto result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
}
};

14
level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_1.cpp
View File

@@ -213,7 +213,7 @@ HWTEST2_P(L0DebuggerParameterizedTests, givenDebuggerWhenAppendingKernelToComman
std::unique_ptr<L0::CommandList> commandList(L0::CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
commandList->close();
@@ -243,7 +243,7 @@ HWTEST2_P(L0DebuggerParameterizedTests, givenDebuggerWhenAppendingKernelToComman
std::unique_ptr<L0::CommandList> commandList(L0::CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
commandList->close();
@@ -607,7 +607,7 @@ HWTEST2_F(L0DebuggerTest, givenDebuggerEnabledAndL1CachePolicyWBWhenAppendingThe
std::unique_ptr<L0::CommandList> commandList(L0::CommandList::create(productFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue));
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, returnValue);
commandList->close();
@@ -668,7 +668,7 @@ HWTEST2_F(L0DebuggerTest, givenFlushTaskSubmissionAndSharedHeapsEnabledWhenAppen
CmdListKernelLaunchParams launchParams = {};
ze_group_count_t groupCount{1, 1, 1};
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
auto csrHeap = &commandList->csr->getIndirectHeap(NEO::HeapType::SURFACE_STATE, 0);
@@ -682,7 +682,7 @@ HWTEST2_F(L0DebuggerTest, givenFlushTaskSubmissionAndSharedHeapsEnabledWhenAppen
memset(debugSurfaceState, 0, sizeof(*debugSurfaceState));
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
ASSERT_EQ(0u, debugSurfaceState->getSurfaceBaseAddress());
@@ -717,7 +717,7 @@ HWTEST2_F(L0DebuggerTest, givenImmediateFlushTaskWhenAppendingKernelUsingNewHeap
CmdListKernelLaunchParams launchParams = {};
ze_group_count_t groupCount{1, 1, 1};
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
auto csrHeap = &commandList->csr->getIndirectHeap(NEO::HeapType::SURFACE_STATE, 0);
@@ -731,7 +731,7 @@ HWTEST2_F(L0DebuggerTest, givenImmediateFlushTaskWhenAppendingKernelUsingNewHeap
memset(debugSurfaceState, 0, sizeof(*debugSurfaceState));
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
returnValue = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, returnValue);
ASSERT_EQ(0u, debugSurfaceState->getSurfaceBaseAddress());

12
level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_2.cpp
View File

@@ -133,7 +133,7 @@ HWTEST2_F(singleAddressSpaceModeTest, givenImmediateCommandListWhenExecutingWith
csr.lastFlushedCommandStream = nullptr;
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_NE(nullptr, csr.lastFlushedCommandStream);
@@ -179,7 +179,7 @@ HWTEST2_P(L0DebuggerWithBlitterTest, givenImmediateCommandListWhenExecutingWithF
csr.lastFlushedCommandStream = nullptr;
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_NE(nullptr, csr.lastFlushedCommandStream);
@@ -224,7 +224,7 @@ HWTEST2_P(L0DebuggerWithBlitterTest, givenImmediateFlushTaskWhenExecutingKernelT
csr.lastFlushedCommandStream = nullptr;
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_NE(nullptr, csr.lastFlushedCommandStream);
@@ -266,7 +266,7 @@ HWTEST_P(L0DebuggerWithBlitterTest, givenInternalUsageImmediateCommandListWhenEx
csr.storeMakeResidentAllocations = true;
CmdListKernelLaunchParams launchParams = {};
auto result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, launchParams, false);
auto result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, nullptr, 0, nullptr, launchParams, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
auto sbaBuffer = device->getL0Debugger()->getSbaTrackingBuffer(commandList->csr->getOsContext().getContextId());
@@ -298,7 +298,7 @@ HWTEST2_P(L0DebuggerWithBlitterTest, givenUseCsrImmediateSubmissionEnabledForImm
ze_group_count_t groupCount{1, 1, 1};
auto commandList = CommandList::createImmediate(productFamily, device, &queueDesc, false, NEO::EngineGroupType::RenderCompute, returnValue);
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, false);
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), groupCount, nullptr, 0, nullptr, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
commandList->destroy();
@@ -315,7 +315,7 @@ HWTEST2_P(L0DebuggerWithBlitterTest, givenUseCsrImmediateSubmissionDisabledForIm
ze_group_count_t groupCount{1, 1, 1};
auto commandList = CommandList::createImmediate(productFamily, device, &queueDesc, false, NEO::EngineGroupType::RenderCompute, returnValue);
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), &groupCount, nullptr, 0, nullptr, false);
auto result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), groupCount, nullptr, 0, nullptr, false);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
commandList->destroy();

22
level_zero/core/test/unit_tests/xe_hpc_core/test_cmdlist_xe_hpc_core.cpp
View File

@@ -48,7 +48,7 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpcCore, givenKernelUsingSyncBufferWhen
kernelAttributes.numGrfRequired = GrfConfig::DefaultGrfNumber;
CmdListKernelLaunchParams launchParams = {};
launchParams.isCooperative = true;
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
{
@@ -56,11 +56,11 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpcCore, givenKernelUsingSyncBufferWhen
VariableBackup<unsigned short> hwRevId{&hwInfo.platform.usRevId};
engineGroupType = EngineGroupType::RenderCompute;
hwRevId = productHelper.getHwRevIdFromStepping(REVISION_B, hwInfo);
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result);
ze_group_count_t groupCount1{1, 1, 1};
result = pCommandList->appendLaunchKernelWithParams(&kernel, &groupCount1, nullptr, launchParams);
result = pCommandList->appendLaunchKernelWithParams(&kernel, groupCount1, nullptr, launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
}
}
@@ -414,7 +414,7 @@ HWTEST2_F(CommandListStatePrefetchXeHpcCore, givenAppendMemoryPrefetchForKmdMigr
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_FALSE(memoryManager->setMemPrefetchCalled);
@@ -466,7 +466,7 @@ HWTEST2_F(CommandListStatePrefetchXeHpcCore, givenAppendMemoryPrefetchForKmdMigr
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
EXPECT_TRUE(memoryManager->setMemPrefetchCalled);
@@ -526,7 +526,7 @@ HWTEST2_F(CommandListStatePrefetchXeHpcCore, givenAppendMemoryPrefetchForKmdMigr
ze_group_count_t groupCount{1, 1, 1};
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel->toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel->toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
commandList->close();
@@ -773,7 +773,7 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpcCore,
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernelWithParams(&kernel, &groupCount, event.get(), launchParams);
result = commandList->appendLaunchKernelWithParams(&kernel, groupCount, event.get(), launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList commands;
@@ -841,7 +841,7 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpcCore,
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernelWithParams(&kernel, &groupCount, event.get(), launchParams);
result = commandList->appendLaunchKernelWithParams(&kernel, groupCount, event.get(), launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList commands;
@@ -969,7 +969,7 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpcCore,
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernelWithParams(&kernel, &groupCount, event.get(), launchParams);
result = commandList->appendLaunchKernelWithParams(&kernel, groupCount, event.get(), launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList commands;
@@ -1039,7 +1039,7 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpcCore,
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernelWithParams(&kernel, &groupCount, event.get(), launchParams);
result = commandList->appendLaunchKernelWithParams(&kernel, groupCount, event.get(), launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList commands;
@@ -1107,7 +1107,7 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpcCore,
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernelWithParams(&kernel, &groupCount, event.get(), launchParams);
result = commandList->appendLaunchKernelWithParams(&kernel, groupCount, event.get(), launchParams);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
GenCmdList commands;

29
level_zero/core/test/unit_tests/xe_hpg_core/dg2/test_cmdlist_dg2.cpp
View File

@@ -80,7 +80,7 @@ HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForRegu
kernel.groupSize[0] = 7;
kernel.groupSize[1] = 1;
kernel.groupSize[2] = 1;
commandList->updateStreamPropertiesForRegularCommandLists(kernel, false, &launchKernelArgs, false);
commandList->updateStreamPropertiesForRegularCommandLists(kernel, false, launchKernelArgs, false);
EXPECT_TRUE(commandList->finalStreamState.frontEndState.disableEUFusion.value);
}
HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForRegularCommandListsCalledAndLwsIsNonOddThenFusedEuIsNotDisabled, IsDG2) {
@@ -97,7 +97,7 @@ HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForRegu
kernel.groupSize[0] = 8;
kernel.groupSize[1] = 1;
kernel.groupSize[2] = 1;
commandList->updateStreamPropertiesForRegularCommandLists(kernel, false, &launchKernelArgs, false);
commandList->updateStreamPropertiesForRegularCommandLists(kernel, false, launchKernelArgs, false);
EXPECT_FALSE(commandList->finalStreamState.frontEndState.disableEUFusion.value);
}
@@ -115,7 +115,7 @@ HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForRegu
kernel.groupSize[0] = 1;
kernel.groupSize[1] = 1;
kernel.groupSize[2] = 1;
commandList->updateStreamPropertiesForRegularCommandLists(kernel, false, &launchKernelArgs, true);
commandList->updateStreamPropertiesForRegularCommandLists(kernel, false, launchKernelArgs, true);
EXPECT_TRUE(commandList->finalStreamState.frontEndState.disableEUFusion.value);
}
@@ -133,7 +133,7 @@ HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForFlus
kernel.groupSize[0] = 7;
kernel.groupSize[1] = 1;
kernel.groupSize[2] = 1;
commandList->updateStreamPropertiesForFlushTaskDispatchFlags(kernel, false, &launchKernelArgs, false);
commandList->updateStreamPropertiesForFlushTaskDispatchFlags(kernel, false, launchKernelArgs, false);
EXPECT_TRUE(commandList->requiredStreamState.frontEndState.disableEUFusion.value);
}
HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForFlushTaskDispatchFlagsCalledAndLwsIsNonOddThenFusedEuIsNotDisabled, IsDG2) {
@@ -150,7 +150,7 @@ HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForFlus
kernel.groupSize[0] = 8;
kernel.groupSize[1] = 1;
kernel.groupSize[2] = 1;
commandList->updateStreamPropertiesForFlushTaskDispatchFlags(kernel, false, &launchKernelArgs, false);
commandList->updateStreamPropertiesForFlushTaskDispatchFlags(kernel, false, launchKernelArgs, false);
EXPECT_FALSE(commandList->requiredStreamState.frontEndState.disableEUFusion.value);
}
@@ -168,26 +168,9 @@ HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForFlus
kernel.groupSize[0] = 1;
kernel.groupSize[1] = 1;
kernel.groupSize[2] = 1;
commandList->updateStreamPropertiesForFlushTaskDispatchFlags(kernel, false, &launchKernelArgs, true);
commandList->updateStreamPropertiesForFlushTaskDispatchFlags(kernel, false, launchKernelArgs, true);
EXPECT_TRUE(commandList->requiredStreamState.frontEndState.disableEUFusion.value);
}
HWTEST2_F(CommandListTests, GivenKernelWithDpasWhenUpdateStreamPropertiesForFlushTaskDispatchFlagsCalledAndGroupCountIsNullptrThenFusedEuIsNotDisabled, IsDG2) {
Mock<::L0::KernelImp> kernel;
auto pMockModule = std::unique_ptr<Module>(new Mock<Module>(device, nullptr));
kernel.module = pMockModule.get();
auto commandList = std::make_unique<WhiteBox<::L0::CommandListCoreFamily<gfxCoreFamily>>>();
auto result = commandList->initialize(device, NEO::EngineGroupType::Compute, 0u);
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.flags.usesSystolicPipelineSelectMode = true;
kernel.groupSize[0] = 1;
kernel.groupSize[1] = 1;
kernel.groupSize[2] = 1;
commandList->updateStreamPropertiesForFlushTaskDispatchFlags(kernel, false, nullptr, true);
EXPECT_FALSE(commandList->requiredStreamState.frontEndState.disableEUFusion.value);
}
} // namespace ult
} // namespace L0

10
level_zero/core/test/unit_tests/xe_hpg_core/test_cmdlist_xe_hpg_core.cpp
View File

@@ -353,7 +353,7 @@ HWTEST2_F(CommandListCreate, GivenComputeModePropertiesWhenUpdateStreamPropertie
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x100;
const ze_group_count_t launchKernelArgs = {};
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
if (commandList->stateComputeModeTracking) {
EXPECT_FALSE(commandList->finalStreamState.stateComputeMode.isCoherencyRequired.isDirty);
if (productHelper.isGrfNumReportedWithScm()) {
@@ -366,7 +366,7 @@ HWTEST2_F(CommandListCreate, GivenComputeModePropertiesWhenUpdateStreamPropertie
EXPECT_TRUE(commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
}
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x80;
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_TRUE(commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
EXPECT_FALSE(commandList->finalStreamState.stateComputeMode.isCoherencyRequired.isDirty);
}
@@ -391,7 +391,7 @@ HWTEST2_F(CommandListCreate, GivenComputeModePropertiesWhenUpdateStreamPropertie
ASSERT_EQ(ZE_RESULT_SUCCESS, result);
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x100;
const ze_group_count_t launchKernelArgs = {};
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
if (commandList->stateComputeModeTracking) {
EXPECT_FALSE(commandList->finalStreamState.stateComputeMode.isCoherencyRequired.isDirty);
EXPECT_FALSE(commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
@@ -400,7 +400,7 @@ HWTEST2_F(CommandListCreate, GivenComputeModePropertiesWhenUpdateStreamPropertie
EXPECT_TRUE(commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
}
const_cast<NEO::KernelDescriptor *>(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x80;
commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false);
commandList->updateStreamProperties(kernel, false, launchKernelArgs, false);
EXPECT_TRUE(commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty);
EXPECT_FALSE(commandList->finalStreamState.stateComputeMode.isCoherencyRequired.isDirty);
}
@@ -441,7 +441,7 @@ HWTEST2_F(CommandListAppendLaunchKernelXeHpgCore, givenEventWhenAppendKernelIsCa
auto usedSpaceBefore = commandList->getCmdContainer().getCommandStream()->getUsed();
CmdListKernelLaunchParams launchParams = {};
result = commandList->appendLaunchKernel(kernel.toHandle(), &groupCount, event->toHandle(), 0, nullptr, launchParams, false);
result = commandList->appendLaunchKernel(kernel.toHandle(), groupCount, event->toHandle(), 0, nullptr, launchParams, false);
EXPECT_EQ(ZE_RESULT_SUCCESS, result);
auto usedSpaceAfter = commandList->getCmdContainer().getCommandStream()->getUsed();