From cb730d11f42774de2a3baeef87a07a7ba9f8cc7f Mon Sep 17 00:00:00 2001 From: Mateusz Jablonski Date: Wed, 27 Sep 2023 16:27:47 +0000 Subject: [PATCH] refactor: pass ze group count as a reference, not a pointer Signed-off-by: Mateusz Jablonski --- .../api/core/ze_module_api_entrypoints.h | 8 +- level_zero/core/source/cmdlist/cmdlist.h | 8 +- level_zero/core/source/cmdlist/cmdlist_hw.h | 20 +- level_zero/core/source/cmdlist/cmdlist_hw.inl | 67 +++--- .../source/cmdlist/cmdlist_hw_immediate.h | 6 +- .../source/cmdlist/cmdlist_hw_immediate.inl | 6 +- .../cmdlist/cmdlist_hw_skl_to_tgllp.inl | 18 +- .../cmdlist/cmdlist_hw_xehp_and_later.inl | 18 +- .../unit_tests/fixtures/cmdlist_fixture.h | 4 +- .../unit_tests/fixtures/cmdlist_fixture.inl | 62 ++--- .../gen12lp/test_cmdlist_gen12lp.cpp | 2 +- ...test_cmdlist_append_launch_kernel_gen9.cpp | 2 +- .../gen9/test_cmdlist_create_gen9.cpp | 8 +- .../core/test/unit_tests/mocks/mock_cmdlist.h | 18 +- .../linux/cmdqueue_linux_tests.cpp | 2 +- .../unit_tests/sources/assert/test_assert.cpp | 12 +- .../sources/cmdlist/test_cmdlist_1.cpp | 42 ++-- .../sources/cmdlist/test_cmdlist_2.cpp | 38 ++-- .../sources/cmdlist/test_cmdlist_3.cpp | 20 +- .../sources/cmdlist/test_cmdlist_4.cpp | 8 +- .../sources/cmdlist/test_cmdlist_5.cpp | 70 +++--- .../sources/cmdlist/test_cmdlist_6.cpp | 42 ++-- .../sources/cmdlist/test_cmdlist_7.cpp | 81 +++---- .../test_cmdlist_append_launch_kernel_1.cpp | 66 +++--- .../test_cmdlist_append_launch_kernel_2.cpp | 44 ++-- .../test_cmdlist_append_launch_kernel_3.cpp | 212 +++++++++--------- .../test_cmdlist_append_wait_on_events.cpp | 6 +- .../cmdlist/test_cmdlist_xehp_and_later.cpp | 82 +++---- .../sources/cmdqueue/test_cmdqueue_1.cpp | 6 +- .../sources/cmdqueue/test_cmdqueue_2.cpp | 2 +- .../sources/cmdqueue/test_cmdqueue_3.cpp | 40 ++-- .../cmdqueue/test_cmdqueue_debugger.cpp | 2 +- .../test_cmdqueue_enqueue_cmdlist.cpp | 4 +- .../test_cmdqueue_enqueue_cmdlist_2.cpp | 6 +- .../sources/debugger/test_l0_debugger_1.cpp | 14 +- .../sources/debugger/test_l0_debugger_2.cpp | 12 +- .../xe_hpc_core/test_cmdlist_xe_hpc_core.cpp | 22 +- .../xe_hpg_core/dg2/test_cmdlist_dg2.cpp | 29 +-- .../xe_hpg_core/test_cmdlist_xe_hpg_core.cpp | 10 +- 39 files changed, 551 insertions(+), 568 deletions(-) diff --git a/level_zero/api/core/ze_module_api_entrypoints.h b/level_zero/api/core/ze_module_api_entrypoints.h index 7160a76e2e..40a5948061 100644 --- a/level_zero/api/core/ze_module_api_entrypoints.h +++ b/level_zero/api/core/ze_module_api_entrypoints.h @@ -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( diff --git a/level_zero/core/source/cmdlist/cmdlist.h b/level_zero/core/source/cmdlist/cmdlist.h index 4ab81ae8be..aef3bd0326 100644 --- a/level_zero/core/source/cmdlist/cmdlist.h +++ b/level_zero/core/source/cmdlist/cmdlist.h @@ -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; diff --git a/level_zero/core/source/cmdlist/cmdlist_hw.h b/level_zero/core/source/cmdlist/cmdlist_hw.h index f586faab8c..64a0a62de0 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw.h +++ b/level_zero/core/source/cmdlist/cmdlist_hw.h @@ -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); diff --git a/level_zero/core/source/cmdlist/cmdlist_hw.inl b/level_zero/core/source/cmdlist/cmdlist_hw.inl index 401c50ba37..b596708d27 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw.inl +++ b/level_zero/core/source/cmdlist/cmdlist_hw.inl @@ -321,7 +321,7 @@ void CommandListCoreFamily::programL3(bool isSLMused) {} template ze_result_t CommandListCoreFamily::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::appendLaunchKernel(ze_kernel_h template ze_result_t CommandListCoreFamily::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::appendLaunchCooperativeKernel( template ze_result_t CommandListCoreFamily::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 ze_result_t CommandListCoreFamily::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::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::appendLaunchMultipleKernelsInd NEO::EncodeMathMMIO::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::appendImageCopyFromMemory(ze_i CmdListKernelLaunchParams launchParams = {}; launchParams.isBuiltInKernel = true; - auto status = CommandListCoreFamily::appendLaunchKernel(builtinKernel->toHandle(), &kernelArgs, + auto status = CommandListCoreFamily::appendLaunchKernel(builtinKernel->toHandle(), kernelArgs, event, numWaitEvents, phWaitEvents, launchParams, relaxedOrderingDispatch); addToMappedEventList(Event::fromHandle(hEvent)); @@ -867,7 +866,7 @@ ze_result_t CommandListCoreFamily::appendImageCopyToMemory(void * launchParams.isDestinationAllocationInSystemMemory = (dstAllocationType == NEO::AllocationType::BUFFER_HOST_MEMORY) || (dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR); - ret = CommandListCoreFamily::appendLaunchKernel(builtinKernel->toHandle(), &kernelArgs, + ret = CommandListCoreFamily::appendLaunchKernel(builtinKernel->toHandle(), kernelArgs, event, numWaitEvents, phWaitEvents, launchParams, relaxedOrderingDispatch); addToMappedEventList(event); @@ -1012,7 +1011,7 @@ ze_result_t CommandListCoreFamily::appendImageCopyRegion(ze_image CmdListKernelLaunchParams launchParams = {}; launchParams.isBuiltInKernel = true; - auto status = CommandListCoreFamily::appendLaunchKernel(kernel->toHandle(), &kernelArgs, + auto status = CommandListCoreFamily::appendLaunchKernel(kernel->toHandle(), kernelArgs, event, numWaitEvents, phWaitEvents, launchParams, relaxedOrderingDispatch); addToMappedEventList(event); @@ -1149,7 +1148,7 @@ ze_result_t CommandListCoreFamily::appendMemoryCopyKernelWithGA(v (dstAllocationType == NEO::AllocationType::SVM_CPU) || (dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR); - return CommandListCoreFamily::appendLaunchKernelSplit(builtinKernel, &dispatchKernelArgs, signalEvent, launchParams); + return CommandListCoreFamily::appendLaunchKernelSplit(builtinKernel, dispatchKernelArgs, signalEvent, launchParams); } template @@ -1647,7 +1646,7 @@ ze_result_t CommandListCoreFamily::appendMemoryCopyKernel3d(Align launchParams.isDestinationAllocationInSystemMemory = (dstAllocationType == NEO::AllocationType::BUFFER_HOST_MEMORY) || (dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR); - return CommandListCoreFamily::appendLaunchKernel(builtinKernel->toHandle(), &dispatchKernelArgs, signalEvent, numWaitEvents, + return CommandListCoreFamily::appendLaunchKernel(builtinKernel->toHandle(), dispatchKernelArgs, signalEvent, numWaitEvents, phWaitEvents, launchParams, relaxedOrderingDispatch); } @@ -1715,7 +1714,7 @@ ze_result_t CommandListCoreFamily::appendMemoryCopyKernel2d(Align (dstAllocationType == NEO::AllocationType::BUFFER_HOST_MEMORY) || (dstAllocationType == NEO::AllocationType::EXTERNAL_HOST_PTR); return CommandListCoreFamily::appendLaunchKernel(builtinKernel->toHandle(), - &dispatchKernelArgs, signalEvent, + dispatchKernelArgs, signalEvent, numWaitEvents, phWaitEvents, launchParams, relaxedOrderingDispatch); @@ -1749,7 +1748,7 @@ ze_result_t CommandListCoreFamily::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::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::appendMemoryFill(void *ptr, builtinKernel->setArgumentValue(3, sizeof(fillArguments.patternSizeInEls), &fillArguments.patternSizeInEls); ze_group_count_t dispatchKernelArgs{static_cast(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::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::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::appendSignalInOrderDependencyCounter( template ze_result_t CommandListCoreFamily::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::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::prepareIndirectParams(const ze } template -void CommandListCoreFamily::updateStreamProperties(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect) { +void CommandListCoreFamily::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::updateStreamProperties(Kernel &kernel } template -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(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 -void CommandListCoreFamily::updateStreamPropertiesForFlushTaskDispatchFlags(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect) { +void CommandListCoreFamily::updateStreamPropertiesForFlushTaskDispatchFlags(Kernel &kernel, bool isCooperative, const ze_group_count_t &threadGroupDimensions, bool isIndirect) { auto &kernelAttributes = kernel.getKernelDescriptor().kernelAttributes; bool fusedEuDisabled = getFusedEuDisabled(kernel, this->device, threadGroupDimensions, isIndirect); @@ -2785,7 +2782,7 @@ void CommandListCoreFamily::updateStreamPropertiesForFlushTaskDis } template -void CommandListCoreFamily::updateStreamPropertiesForRegularCommandLists(Kernel &kernel, bool isCooperative, const ze_group_count_t *threadGroupDimensions, bool isIndirect) { +void CommandListCoreFamily::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; diff --git a/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h b/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h index fdf38e8b09..2cc86ae874 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h +++ b/level_zero/core/source/cmdlist/cmdlist_hw_immediate.h @@ -52,13 +52,13 @@ struct CommandListCoreFamilyImmediate : public CommandListCoreFamily::skipInOrderNonWalkerSignalin template ze_result_t CommandListCoreFamilyImmediate::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::handleInOrderNonWalkerSignal template ze_result_t CommandListCoreFamilyImmediate::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::appendMemoryRangesBar template ze_result_t CommandListCoreFamilyImmediate::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) { diff --git a/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl b/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl index 0a7e1b7313..dcf8077e7d 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl +++ b/level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl @@ -47,7 +47,7 @@ bool CommandListCoreFamily::isInOrderNonWalkerSignalingRequired(c template ze_result_t CommandListCoreFamily::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::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::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::appendLaunchKernelWithParams(K kernel, // dispatchInterface ssh, // surfaceStateHeap dsh, // dynamicStateHeap - reinterpret_cast(threadGroupDimensions), // threadGroupDimensions + reinterpret_cast(&threadGroupDimensions), // threadGroupDimensions nullptr, // outWalkerPtr &additionalCommands, // additionalCommands commandListPreemptionMode, // preemptionMode @@ -300,7 +300,7 @@ inline size_t CommandListCoreFamily::estimateBufferSizeMultiTileB template ze_result_t CommandListCoreFamily::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); diff --git a/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl b/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl index 3b4bf784cd..6e8a59b70c 100644 --- a/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl +++ b/level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl @@ -82,7 +82,7 @@ bool CommandListCoreFamily::isInOrderNonWalkerSignalingRequired(c } template -ze_result_t CommandListCoreFamily::appendLaunchKernelWithParams(Kernel *kernel, const ze_group_count_t *threadGroupDimensions, Event *event, +ze_result_t CommandListCoreFamily::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::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(commandListPerThreadScratchSize, kernelDescriptor.kernelAttributes.perThreadScratchSize[0]); @@ -165,13 +165,13 @@ ze_result_t CommandListCoreFamily::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::appendLaunchKernelWithParams(K kernel, // dispatchInterface ssh, // surfaceStateHeap dsh, // dynamicStateHeap - reinterpret_cast(threadGroupDimensions), // threadGroupDimensions + reinterpret_cast(&threadGroupDimensions), // threadGroupDimensions nullptr, // outWalkerPtr &additionalCommands, // additionalCommands kernelPreemptionMode, // preemptionMode @@ -461,7 +461,7 @@ inline size_t CommandListCoreFamily::estimateBufferSizeMultiTileB template ze_result_t CommandListCoreFamily::appendLaunchKernelSplit(Kernel *kernel, - const ze_group_count_t *threadGroupDimensions, + const ze_group_count_t &threadGroupDimensions, Event *event, const CmdListKernelLaunchParams &launchParams) { if (event) { diff --git a/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.h b/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.h index 213c4ecfc5..3cfb880b0b 100644 --- a/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.h +++ b/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.h @@ -219,14 +219,14 @@ class AppendFillFixture : public DeviceFixture { MockCommandList() : WhiteBox<::L0::CommandListCoreFamily>() {} 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++; diff --git a/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.inl b/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.inl index df4e64cca8..82c5631c86 100644 --- a/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.inl +++ b/level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.inl @@ -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); diff --git a/level_zero/core/test/unit_tests/gen12lp/test_cmdlist_gen12lp.cpp b/level_zero/core/test/unit_tests/gen12lp/test_cmdlist_gen12lp.cpp index 6a7056dd4e..00c83cb7ed 100644 --- a/level_zero/core/test/unit_tests/gen12lp/test_cmdlist_gen12lp.cpp +++ b/level_zero/core/test/unit_tests/gen12lp/test_cmdlist_gen12lp.cpp @@ -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(); diff --git a/level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp b/level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp index 64a4b8f042..4a29c06eae 100644 --- a/level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp +++ b/level_zero/core/test/unit_tests/gen9/test_cmdlist_append_launch_kernel_gen9.cpp @@ -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(); diff --git a/level_zero/core/test/unit_tests/gen9/test_cmdlist_create_gen9.cpp b/level_zero/core/test/unit_tests/gen9/test_cmdlist_create_gen9.cpp index a85f5d2648..6eeab70fc0 100644 --- a/level_zero/core/test/unit_tests/gen9/test_cmdlist_create_gen9.cpp +++ b/level_zero/core/test/unit_tests/gen9/test_cmdlist_create_gen9.cpp @@ -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(); diff --git a/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h b/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h index d7b143630a..92177c261b 100644 --- a/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h +++ b/level_zero/core/test/unit_tests/mocks/mock_cmdlist.h @@ -105,7 +105,7 @@ struct WhiteBox<::L0::CommandListCoreFamily> WhiteBox() : ::L0::CommandListCoreFamily(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> 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> } 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; diff --git a/level_zero/core/test/unit_tests/os_interface/linux/cmdqueue_linux_tests.cpp b/level_zero/core/test/unit_tests/os_interface/linux/cmdqueue_linux_tests.cpp index 50653a1e17..c082bf81cf 100644 --- a/level_zero/core/test/unit_tests/os_interface/linux/cmdqueue_linux_tests.cpp +++ b/level_zero/core/test/unit_tests/os_interface/linux/cmdqueue_linux_tests.cpp @@ -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()}; diff --git a/level_zero/core/test/unit_tests/sources/assert/test_assert.cpp b/level_zero/core/test/unit_tests/sources/assert/test_assert.cpp index 272998b65c..9185e90354 100644 --- a/level_zero/core/test/unit_tests/sources/assert/test_assert.cpp +++ b/level_zero/core/test/unit_tests/sources/assert/test_assert.cpp @@ -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(); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp index 2de6b84af6..4dae26b55c 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_1.cpp @@ -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(&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(Event::fromHandle(event2)->getHostAddress()) = static_cast(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( diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp index 2187513765..08953cdc31 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_2.cpp @@ -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(); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_3.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_3.cpp index 31df6c0f8b..945b30623b 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_3.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_3.cpp @@ -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) { diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp index e997d087f9..063539e3af 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_4.cpp @@ -331,7 +331,7 @@ HWTEST_F(CommandListImmediateFlushTaskComputeTests, givenUseCsrImmediateSubmissi std::unique_ptr 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(); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_5.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_5.cpp index 625f7f4146..4df90b18c8 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_5.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_5.cpp @@ -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(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(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); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_6.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_6.cpp index 5825d70cc1..b3eec8fe98 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_6.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_6.cpp @@ -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().bindless = undefined; ptrArg.as().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); } diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_7.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_7.cpp index d2bfc490e5..1e8fb3b3f4 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_7.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_7.cpp @@ -264,7 +264,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenSignalEventWhenAppendLaunchCoopera auto commandList = std::make_unique>>(); 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(&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(alloc), event->toHandle(), 0, nullptr, false); + result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), *static_cast(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(&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(&kernel.getKernelDescriptor())->kernelAttributes.numGrfRequired = 0x80; - commandList->updateStreamProperties(kernel, false, &launchKernelArgs, false); + commandList->updateStreamProperties(kernel, false, launchKernelArgs, false); if constexpr (TestTraits::largeGrfModeInStateComputeModeSupported) { EXPECT_EQ(productHelper.isGrfNumReportedWithScm(), commandList->finalStreamState.stateComputeMode.largeGrfMode.isDirty); } @@ -403,7 +404,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, const_cast(&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(&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(&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(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(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 ®ularCmdListStream = *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 class MockCommandListHwKernelSplit : public WhiteBox<::L0::CommandListCoreFamily> { 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>>(); 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>>(); 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>>(); 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); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_1.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_1.cpp index 2ede64ecee..3f6f9d078e 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_1.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_1.cpp @@ -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 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 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 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(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(alloc), + *static_cast(alloc), nullptr, 0, nullptr, false); EXPECT_EQ(result, ZE_RESULT_SUCCESS); kernel.groupSize[2] = 2; result = commandList->appendLaunchKernelIndirect(kernel.toHandle(), - static_cast(alloc), + *static_cast(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 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); } } diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_2.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_2.cpp index f350d3d602..9b84c5a49b 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_2.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_2.cpp @@ -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(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>>(); 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(); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp index beed7b6786..5bbe8f5ee0 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_launch_kernel_3.cpp @@ -62,7 +62,7 @@ HWCMDTEST_F(IGFX_GEN8_CORE, CommandListAppendLaunchKernel, givenFunctionWhenBind ze_result_t returnValue; std::unique_ptr 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::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(&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(&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::create(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::create(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>>(); 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>>(); 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>>(); 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>>(); 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>>(); 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 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>>(); 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>>(); 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 *>(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(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(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(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(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 *>(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 *>(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(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(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(*iterator); @@ -2206,9 +2208,9 @@ HWTEST2_F(InOrderCmdListTests, givenImmediateEventWhenWaitingFromRegularCmdListT auto eventPool = createEvents(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(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(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(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(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(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(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(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(alloc), + *static_cast(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(); diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_wait_on_events.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_wait_on_events.cpp index 1da7aa40fb..885322fc2d 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_wait_on_events.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_append_wait_on_events.cpp @@ -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); } diff --git a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_xehp_and_later.cpp b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_xehp_and_later.cpp index 16c00199f1..35cb82b2ef 100644 --- a/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_xehp_and_later.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdlist/test_cmdlist_xehp_and_later.cpp @@ -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(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); diff --git a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_1.cpp b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_1.cpp index 7749a2d32c..68853ed992 100644 --- a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_1.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_1.cpp @@ -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()}; diff --git a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_2.cpp b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_2.cpp index 582deabdd0..9115bbadcc 100644 --- a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_2.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_2.cpp @@ -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(device); commandList->close(); diff --git a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_3.cpp b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_3.cpp index 101e12f1ed..4279d5d411 100644 --- a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_3.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_3.cpp @@ -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>>(); 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>>(); 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>>(); 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>>(); 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>>(); 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>>(); @@ -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>>(); 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()}; diff --git a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_debugger.cpp b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_debugger.cpp index 8c7e92c3c1..ab8e2503c1 100644 --- a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_debugger.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_debugger.cpp @@ -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(); diff --git a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist.cpp b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist.cpp index 7d55df02e4..70db773e84 100644 --- a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist.cpp @@ -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); diff --git a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist_2.cpp b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist_2.cpp index f6fed827cc..c77154a011 100644 --- a/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist_2.cpp +++ b/level_zero/core/test/unit_tests/sources/cmdqueue/test_cmdqueue_enqueue_cmdlist_2.cpp @@ -380,7 +380,7 @@ struct PauseOnGpuFixture : public Test { } 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); } }; diff --git a/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_1.cpp b/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_1.cpp index 378a931e86..e04e8454b7 100644 --- a/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_1.cpp +++ b/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_1.cpp @@ -213,7 +213,7 @@ HWTEST2_P(L0DebuggerParameterizedTests, givenDebuggerWhenAppendingKernelToComman std::unique_ptr 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 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 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()); diff --git a/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_2.cpp b/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_2.cpp index 00f8090d09..5dec77cfe8 100644 --- a/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_2.cpp +++ b/level_zero/core/test/unit_tests/sources/debugger/test_l0_debugger_2.cpp @@ -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(); diff --git a/level_zero/core/test/unit_tests/xe_hpc_core/test_cmdlist_xe_hpc_core.cpp b/level_zero/core/test/unit_tests/xe_hpc_core/test_cmdlist_xe_hpc_core.cpp index bbae083f04..85414ed0bf 100644 --- a/level_zero/core/test/unit_tests/xe_hpc_core/test_cmdlist_xe_hpc_core.cpp +++ b/level_zero/core/test/unit_tests/xe_hpc_core/test_cmdlist_xe_hpc_core.cpp @@ -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 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; diff --git a/level_zero/core/test/unit_tests/xe_hpg_core/dg2/test_cmdlist_dg2.cpp b/level_zero/core/test/unit_tests/xe_hpg_core/dg2/test_cmdlist_dg2.cpp index fb20f92539..4be255fd2a 100644 --- a/level_zero/core/test/unit_tests/xe_hpg_core/dg2/test_cmdlist_dg2.cpp +++ b/level_zero/core/test/unit_tests/xe_hpg_core/dg2/test_cmdlist_dg2.cpp @@ -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(new Mock(device, nullptr)); - kernel.module = pMockModule.get(); - - auto commandList = std::make_unique>>(); - auto result = commandList->initialize(device, NEO::EngineGroupType::Compute, 0u); - ASSERT_EQ(ZE_RESULT_SUCCESS, result); - - const_cast(&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 diff --git a/level_zero/core/test/unit_tests/xe_hpg_core/test_cmdlist_xe_hpg_core.cpp b/level_zero/core/test/unit_tests/xe_hpg_core/test_cmdlist_xe_hpg_core.cpp index 0bc7d81f1d..4182158b72 100644 --- a/level_zero/core/test/unit_tests/xe_hpg_core/test_cmdlist_xe_hpg_core.cpp +++ b/level_zero/core/test/unit_tests/xe_hpg_core/test_cmdlist_xe_hpg_core.cpp @@ -353,7 +353,7 @@ HWTEST2_F(CommandListCreate, GivenComputeModePropertiesWhenUpdateStreamPropertie const_cast(&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(&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(&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(&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();