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 165355bb8d..c4f7cce0fe 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 @@ -440,7 +440,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenImmediateCommandListWhenAppendingL ASSERT_EQ(ZE_RESULT_ERROR_INVALID_ARGUMENT, result); } -HWTEST2_F(CommandListAppendLaunchKernel, givenKernelUsingSyncBufferWhenAppendLaunchCooperativeKernelIsCalledThenCorrectValueIsReturned, MatchAny) { +HWTEST2_F(CommandListAppendLaunchKernel, givenKernelUsingSyncBufferWhenAppendLaunchCooperativeKernelIsCalledThenCorrectValueIsReturned, HasDispatchAllSupport) { Mock<::L0::KernelImp> kernel; auto pMockModule = std::unique_ptr(new Mock(device, nullptr)); kernel.module = pMockModule.get(); @@ -716,7 +716,7 @@ HWTEST2_F(CommandListAppendLaunchKernel, givenKernelUsingRegionGroupBarrierWhenA EXPECT_EQ(nullptr, kernel.getRegionGroupBarrierAllocation()); } -HWTEST2_F(CommandListAppendLaunchKernel, whenAppendLaunchCooperativeKernelAndQueryKernelTimestampsToTheSameCmdlistThenFronEndStateIsNotChanged, MatchAny) { +HWTEST2_F(CommandListAppendLaunchKernel, whenAppendLaunchCooperativeKernelAndQueryKernelTimestampsToTheSameCmdlistThenFronEndStateIsNotChanged, HasDispatchAllSupport) { Mock<::L0::KernelImp> kernel; auto pMockModule = std::unique_ptr(new Mock(device, nullptr)); kernel.module = pMockModule.get(); diff --git a/level_zero/core/test/unit_tests/sources/kernel/test_kernel_2.cpp b/level_zero/core/test/unit_tests/sources/kernel/test_kernel_2.cpp index 38c7e55e09..12109ddf75 100644 --- a/level_zero/core/test/unit_tests/sources/kernel/test_kernel_2.cpp +++ b/level_zero/core/test/unit_tests/sources/kernel/test_kernel_2.cpp @@ -446,7 +446,7 @@ TEST(zeKernelGetProperties, WhenGettingKernelPropertiesThenSuccessIsReturned) { using KernelImpSuggestMaxCooperativeGroupCountTests = Test; -HWTEST_F(KernelImpSuggestMaxCooperativeGroupCountTests, GivenNoBarriersOrSlmUsedWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithSimd) { +HWTEST2_F(KernelImpSuggestMaxCooperativeGroupCountTests, GivenNoBarriersOrSlmUsedWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithSimd, IsAtLeastBmg) { auto workGroupSize = lws[0] * lws[1] * lws[2]; auto expected = availableThreadCount / Math::divideAndRoundUp(workGroupSize, simd); EXPECT_EQ(expected, getMaxWorkGroupCount()); @@ -472,13 +472,13 @@ HWTEST_F(KernelImpSuggestMaxCooperativeGroupCountTests, GivenMultiTileWhenCalcul } } -HWTEST_F(KernelImpSuggestMaxCooperativeGroupCountTests, GivenBarriersWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithRegardToBarriersCount) { +HWTEST2_F(KernelImpSuggestMaxCooperativeGroupCountTests, GivenBarriersWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithRegardToBarriersCount, IsAtLeastBmg) { usesBarriers = 1; auto expected = dssCount * (maxBarrierCount / usesBarriers); EXPECT_EQ(expected, getMaxWorkGroupCount()); } -HWTEST_F(KernelImpSuggestMaxCooperativeGroupCountTests, GivenUsedSlmSizeWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithRegardToUsedSlmSize) { +HWTEST2_F(KernelImpSuggestMaxCooperativeGroupCountTests, GivenUsedSlmSizeWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithRegardToUsedSlmSize, IsAtLeastBmg) { usedSlm = 64 * MemoryConstants::kiloByte; auto expected = availableSlm / usedSlm; EXPECT_EQ(expected, getMaxWorkGroupCount()); diff --git a/level_zero/core/test/unit_tests/xe_hpc_core/pvc/test_kernel_pvc.cpp b/level_zero/core/test/unit_tests/xe_hpc_core/pvc/test_kernel_pvc.cpp index e9877bdc6e..96bf2cccf5 100644 --- a/level_zero/core/test/unit_tests/xe_hpc_core/pvc/test_kernel_pvc.cpp +++ b/level_zero/core/test/unit_tests/xe_hpc_core/pvc/test_kernel_pvc.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2023-2024 Intel Corporation + * Copyright (C) 2023-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -13,11 +13,6 @@ #include "level_zero/core/test/unit_tests/fixtures/kernel_max_cooperative_groups_count_fixture.h" namespace L0 { namespace ult { - -HWTEST_EXCLUDE_PRODUCT(KernelImpSuggestMaxCooperativeGroupCountTests, GivenUsedSlmSizeWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithRegardToUsedSlmSize, IGFX_PVC); -HWTEST_EXCLUDE_PRODUCT(KernelImpSuggestMaxCooperativeGroupCountTests, GivenBarriersWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithRegardToBarriersCount, IGFX_PVC); -HWTEST_EXCLUDE_PRODUCT(KernelImpSuggestMaxCooperativeGroupCountTests, GivenNoBarriersOrSlmUsedWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithSimd, IGFX_PVC); - using KernelImpSuggestMaxCooperativeGroupCountTestsPvc = Test; PVCTEST_F(KernelImpSuggestMaxCooperativeGroupCountTestsPvc, GivenNoBarriersOrSlmUsedWhenCalculatingMaxCooperativeGroupCountThenResultIsCalculatedWithSimd) { diff --git a/opencl/test/unit_test/command_queue/sync_buffer_handler_tests.cpp b/opencl/test/unit_test/command_queue/sync_buffer_handler_tests.cpp index 3e01afd060..90b2bc2f17 100644 --- a/opencl/test/unit_test/command_queue/sync_buffer_handler_tests.cpp +++ b/opencl/test/unit_test/command_queue/sync_buffer_handler_tests.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2019-2024 Intel Corporation + * Copyright (C) 2019-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -114,7 +114,7 @@ class SyncBufferHandlerTest : public SyncBufferEnqueueHandlerTest { MockCommandQueue *commandQueue; }; -HWTEST_TEMPLATED_F(SyncBufferHandlerTest, GivenAllocateSyncBufferPatchAndConcurrentKernelWhenEnqueuingKernelThenSyncBufferIsUsed) { +HWTEST2_TEMPLATED_F(SyncBufferHandlerTest, GivenAllocateSyncBufferPatchAndConcurrentKernelWhenEnqueuingKernelThenSyncBufferIsUsed, HasDispatchAllSupport) { patchAllocateSyncBuffer(); enqueueNDCount(); @@ -142,7 +142,7 @@ HWTEST_TEMPLATED_F(SyncBufferHandlerTest, GivenAllocateSyncBufferPatchAndConcurr EXPECT_EQ(2u * minimalSyncBufferSize, syncBufferHandler->usedBufferSize); } -HWTEST_TEMPLATED_F(SyncBufferHandlerTest, GivenConcurrentKernelWithoutAllocateSyncBufferPatchWhenEnqueuingConcurrentKernelThenSyncBufferIsNotCreated) { +HWTEST2_TEMPLATED_F(SyncBufferHandlerTest, GivenConcurrentKernelWithoutAllocateSyncBufferPatchWhenEnqueuingConcurrentKernelThenSyncBufferIsNotCreated, HasDispatchAllSupport) { auto retVal = enqueueNDCount(); EXPECT_EQ(CL_SUCCESS, retVal); EXPECT_EQ(nullptr, getSyncBufferHandler()); @@ -157,7 +157,7 @@ HWTEST_TEMPLATED_F(SyncBufferHandlerTest, GivenDefaultKernelUsingSyncBufferWhenE EXPECT_EQ(nullptr, getSyncBufferHandler()); } -HWTEST_TEMPLATED_F(SyncBufferHandlerTest, GivenConcurrentKernelWithAllocateSyncBufferPatchWhenEnqueuingConcurrentKernelThenSyncBufferIsCreated) { +HWTEST2_TEMPLATED_F(SyncBufferHandlerTest, GivenConcurrentKernelWithAllocateSyncBufferPatchWhenEnqueuingConcurrentKernelThenSyncBufferIsCreated, HasDispatchAllSupport) { patchAllocateSyncBuffer(); auto retVal = enqueueNDCount(); EXPECT_EQ(CL_SUCCESS, retVal); @@ -180,7 +180,7 @@ HWTEST_TEMPLATED_F(SyncBufferHandlerTest, GivenTooHighWorkgroupCountWhenEnqueuin EXPECT_EQ(CL_INVALID_VALUE, retVal); } -HWTEST_TEMPLATED_F(SyncBufferHandlerTest, GivenSyncBufferFullWhenEnqueuingKernelThenNewBufferIsAllocated) { +HWTEST2_TEMPLATED_F(SyncBufferHandlerTest, GivenSyncBufferFullWhenEnqueuingKernelThenNewBufferIsAllocated, HasDispatchAllSupport) { patchAllocateSyncBuffer(); enqueueNDCount(); auto syncBufferHandler = getSyncBufferHandler(); diff --git a/shared/source/helpers/gfx_core_helper_bdw_to_dg2.inl b/shared/source/helpers/gfx_core_helper_bdw_to_dg2.inl index f256944263..e4bf694f3f 100644 --- a/shared/source/helpers/gfx_core_helper_bdw_to_dg2.inl +++ b/shared/source/helpers/gfx_core_helper_bdw_to_dg2.inl @@ -1,5 +1,5 @@ /* - * Copyright (C) 2021-2024 Intel Corporation + * Copyright (C) 2021-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -31,7 +31,11 @@ bool GfxCoreHelperHw::isCooperativeDispatchSupported(const EngineGrou template uint32_t GfxCoreHelperHw::adjustMaxWorkGroupCount(uint32_t maxWorkGroupCount, const EngineGroupType engineGroupType, const RootDeviceEnvironment &rootDeviceEnvironment) const { - return maxWorkGroupCount; + if ((debugManager.flags.ForceTheoreticalMaxWorkGroupCount.get()) || + (debugManager.flags.OverrideMaxWorkGroupCount.get() != -1)) { + return maxWorkGroupCount; + } + return 1u; } template diff --git a/shared/test/common/test_macros/header/common_matchers.h b/shared/test/common/test_macros/header/common_matchers.h index ca584895eb..9b884c915f 100644 --- a/shared/test/common/test_macros/header/common_matchers.h +++ b/shared/test/common/test_macros/header/common_matchers.h @@ -96,6 +96,20 @@ using HasOclocZebinFormatEnforced = IsAnyProducts; +struct HasDispatchAllSupport { + template + static constexpr bool isMatched() { + return IsPVC::isMatched() || IsAtLeastBmg::isMatched(); + } +}; + +struct DoesNotHaveDispatchAllSupport { + template + static constexpr bool isMatched() { + return !IsPVC::isMatched() && IsAtMostArl::isMatched(); + } +}; + struct IsXeLpg { template static constexpr bool isMatched() { diff --git a/shared/test/unit_test/helpers/gfx_core_helper_tests.cpp b/shared/test/unit_test/helpers/gfx_core_helper_tests.cpp index fd2b26f86d..726b7b10be 100644 --- a/shared/test/unit_test/helpers/gfx_core_helper_tests.cpp +++ b/shared/test/unit_test/helpers/gfx_core_helper_tests.cpp @@ -1582,7 +1582,7 @@ HWTEST_F(ProductHelperCommonTest, givenProductHelperWhenCallingIsCalculationForD auto &gfxCoreHelper = getHelper(); EXPECT_FALSE(gfxCoreHelper.isCalculationForDisablingEuFusionWithDpasNeeded(hwInfo)); } -HWTEST_F(GfxCoreHelperTest, GivenCooperativeEngineSupportedAndNotUsedWhenAdjustMaxWorkGroupCountIsCalledThenSmallerValueIsReturned) { +HWTEST2_F(GfxCoreHelperTest, GivenCooperativeEngineSupportedAndNotUsedWhenAdjustMaxWorkGroupCountIsCalledThenSmallerValueIsReturned, IsAtLeastBmg) { MockExecutionEnvironment mockExecutionEnvironment{}; auto &rootDeviceEnvironment = *mockExecutionEnvironment.rootDeviceEnvironments[0]; @@ -1631,6 +1631,24 @@ HWTEST_F(GfxCoreHelperTest, GivenCooperativeEngineSupportedAndNotUsedWhenAdjustM } } +HWTEST2_F(GfxCoreHelperTest, GivenGfxHelperWhenAdjustMaxWorkGroupSizeCalledThenOneReturned, DoesNotHaveDispatchAllSupport) { + MockExecutionEnvironment mockExecutionEnvironment{}; + auto &rootDeviceEnvironment = *mockExecutionEnvironment.rootDeviceEnvironments[0]; + auto &gfxCoreHelper = rootDeviceEnvironment.getHelper(); + auto engineGroupType = EngineGroupType::compute; + EXPECT_EQ(1u, gfxCoreHelper.adjustMaxWorkGroupCount(1024u, engineGroupType, rootDeviceEnvironment)); +} + +HWTEST2_F(GfxCoreHelperTest, GivenGfxHelperWhenForceTheoreticalMaxWorkGroupCountThenAdjustMaxWorkGroupSizeReturnNotOne, DoesNotHaveDispatchAllSupport) { + DebugManagerStateRestore dbgRestore; + debugManager.flags.ForceTheoreticalMaxWorkGroupCount.set(true); + MockExecutionEnvironment mockExecutionEnvironment{}; + auto &rootDeviceEnvironment = *mockExecutionEnvironment.rootDeviceEnvironments[0]; + auto &gfxCoreHelper = rootDeviceEnvironment.getHelper(); + auto engineGroupType = EngineGroupType::compute; + EXPECT_NE(1u, gfxCoreHelper.adjustMaxWorkGroupCount(1024u, engineGroupType, rootDeviceEnvironment)); +} + HWTEST_F(GfxCoreHelperTest, givenNumGrfAndSimdSizeWhenAdjustingMaxWorkGroupSizeThenAlwaysReturnDeviceDefault) { const auto &gfxCoreHelper = getHelper(); const auto &rootDeviceEnvironment = pDevice->getRootDeviceEnvironment(); diff --git a/shared/test/unit_test/helpers/kernel_helpers_tests.cpp b/shared/test/unit_test/helpers/kernel_helpers_tests.cpp index 8929784835..d03196b622 100644 --- a/shared/test/unit_test/helpers/kernel_helpers_tests.cpp +++ b/shared/test/unit_test/helpers/kernel_helpers_tests.cpp @@ -103,7 +103,7 @@ TEST_F(KernelHelperMaxWorkGroupsTests, givenMultipleSubdevicesWenCalculatingMaxW } } -HWTEST2_F(KernelHelperMaxWorkGroupsTests, GivenBarriersWhenCalculatingMaxWorkGroupsCountThenResultIsCalculatedWithRegardToBarriersCount, MatchAny) { +HWTEST2_F(KernelHelperMaxWorkGroupsTests, GivenBarriersWhenCalculatingMaxWorkGroupsCountThenResultIsCalculatedWithRegardToBarriersCount, HasDispatchAllSupport) { NEO::RAIIProductHelperFactory> raii(*rootDeviceEnvironment); raii.mockProductHelper->isCooperativeEngineSupportedValue = false; lws[0] = 1; @@ -121,7 +121,7 @@ HWTEST2_F(KernelHelperMaxWorkGroupsTests, GivenBarriersWhenCalculatingMaxWorkGro EXPECT_EQ(expected, getMaxWorkGroupCount()); } -HWTEST2_F(KernelHelperMaxWorkGroupsTests, GivenUsedSlmSizeWhenCalculatingMaxWorkGroupsCountThenResultIsCalculatedWithRegardToUsedSlmSize, MatchAny) { +HWTEST2_F(KernelHelperMaxWorkGroupsTests, GivenUsedSlmSizeWhenCalculatingMaxWorkGroupsCountThenResultIsCalculatedWithRegardToUsedSlmSize, HasDispatchAllSupport) { NEO::RAIIProductHelperFactory> raii(*rootDeviceEnvironment); raii.mockProductHelper->isCooperativeEngineSupportedValue = false; usedSlm = 0; @@ -157,7 +157,7 @@ HWTEST_F(KernelHelperMaxWorkGroupsTests, givenZeroBarriersAndSlmNotUsedWhenCalcu EXPECT_EQ(raiiFactory.mockGfxCoreHelper->alignThreadGroupCountToDssSizeCalledTimes, 0u); } -TEST_F(KernelHelperMaxWorkGroupsTests, GivenVariousValuesWhenCalculatingMaxWorkGroupsCountThenLowestResultIsAlwaysReturned) { +HWTEST2_F(KernelHelperMaxWorkGroupsTests, GivenVariousValuesWhenCalculatingMaxWorkGroupsCountThenLowestResultIsAlwaysReturned, HasDispatchAllSupport) { auto &helper = rootDeviceEnvironment->getHelper(); engineType = EngineGroupType::cooperativeCompute; diff --git a/shared/test/unit_test/xe_hpc_core/pvc/excludes_xe_hpc_core_pvc.cpp b/shared/test/unit_test/xe_hpc_core/pvc/excludes_xe_hpc_core_pvc.cpp index d9b469abd7..ef236db237 100644 --- a/shared/test/unit_test/xe_hpc_core/pvc/excludes_xe_hpc_core_pvc.cpp +++ b/shared/test/unit_test/xe_hpc_core/pvc/excludes_xe_hpc_core_pvc.cpp @@ -1,5 +1,5 @@ /* - * Copyright (C) 2022-2024 Intel Corporation + * Copyright (C) 2022-2025 Intel Corporation * * SPDX-License-Identifier: MIT * @@ -8,5 +8,4 @@ #include "shared/test/common/test_macros/hw_test_base.h" HWTEST_EXCLUDE_PRODUCT(ProductHelperTest, givenProductHelperWhenAskedIfIsBlitSplitEnqueueWARequiredThenReturnFalse, IGFX_PVC); -HWTEST_EXCLUDE_PRODUCT(BlitTests, GivenCpuAccessToLocalMemoryWhenGettingMaxBlitSizeThenValuesAreOverriden_BlitPlatforms, IGFX_PVC); -HWTEST_EXCLUDE_PRODUCT(GfxCoreHelperTest, GivenCooperativeEngineSupportedAndNotUsedWhenAdjustMaxWorkGroupCountIsCalledThenSmallerValueIsReturned, IGFX_PVC); +HWTEST_EXCLUDE_PRODUCT(BlitTests, GivenCpuAccessToLocalMemoryWhenGettingMaxBlitSizeThenValuesAreOverriden_BlitPlatforms, IGFX_PVC); \ No newline at end of file