refactor: pass extra walker params

Signed-off-by: Bartosz Dunajski <[email protected]>

Author: BartoszDunajski

level_zero/api/core/ze_module_api_entrypoints.h

@@ -103,7 +103,8 @@ ze_result_t zeKernelSuggestGroupSize(
 ze_result_t zeKernelSuggestMaxCooperativeGroupCount(
     ze_kernel_handle_t hKernel,
     uint32_t *totalGroupCount) {
-    return L0::Kernel::fromHandle(hKernel)->suggestMaxCooperativeGroupCount(totalGroupCount, NEO::EngineGroupType::compute, false);
+    *totalGroupCount = L0::Kernel::fromHandle(hKernel)->suggestMaxCooperativeGroupCount(NEO::EngineGroupType::compute, false, false);
+    return ZE_RESULT_SUCCESS;
 }
 
 ze_result_t zeKernelSetArgumentValue(

level_zero/core/source/cmdlist/cmdlist_hw.inl

@@ -2758,12 +2758,9 @@ void CommandListCoreFamily<gfxCoreFamily>::appendSignalInOrderDependencyCounter(
 template <GFXCORE_FAMILY gfxCoreFamily>
 ze_result_t CommandListCoreFamily<gfxCoreFamily>::programSyncBuffer(Kernel &kernel, NEO::Device &device,
                                                                     const ze_group_count_t &threadGroupDimensions) {
-    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);
+    uint32_t maximalNumberOfWorkgroupsAllowed = kernel.suggestMaxCooperativeGroupCount(this->engineGroupType, device.isEngineInstanced(), false);
+
+    size_t requestedNumberOfWorkgroups = (threadGroupDimensions.groupCountX * threadGroupDimensions.groupCountY * threadGroupDimensions.groupCountZ);
     if (requestedNumberOfWorkgroups > maximalNumberOfWorkgroupsAllowed) {
         return ZE_RESULT_ERROR_INVALID_ARGUMENT;
     }

level_zero/core/source/cmdlist/cmdlist_hw_skl_to_tgllp.inl

@@ -191,6 +191,9 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
     std::list<void *> additionalCommands;
 
     updateStreamProperties(*kernel, launchParams.isCooperative, threadGroupDimensions, launchParams.isIndirect);
+
+    auto maxWgCountPerTile = kernel->suggestMaxCooperativeGroupCount(this->engineGroupType, device->getNEODevice()->isEngineInstanced(), true);
+
     NEO::EncodeDispatchKernelArgs dispatchKernelArgs{
         0,                                                      // eventAddress
         static_cast<uint64_t>(Event::STATE_SIGNALED),           // postSyncImmValue
@@ -212,6 +215,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
         launchParams.additionalSizeParam,                       // additionalSizeParam
         0,                                                      // partitionCount
         launchParams.reserveExtraPayloadSpace,                  // reserveExtraPayloadSpace
+        maxWgCountPerTile,                                      // maxWgCountPerTile
         NEO::ThreadArbitrationPolicy::NotPresent,               // defaultPipelinedThreadArbitrationPolicy
         launchParams.isIndirect,                                // isIndirect
         launchParams.isPredicate,                               // isPredicate

level_zero/core/source/cmdlist/cmdlist_hw_xehp_and_later.inl

@@ -345,6 +345,8 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
         }
     }
 
+    auto maxWgCountPerTile = kernel->suggestMaxCooperativeGroupCount(this->engineGroupType, device->getNEODevice()->isEngineInstanced(), true);
+
     NEO::EncodeDispatchKernelArgs dispatchKernelArgs{
         eventAddress,                                           // eventAddress
         static_cast<uint64_t>(Event::STATE_SIGNALED),           // postSyncImmValue
@@ -366,6 +368,7 @@ ze_result_t CommandListCoreFamily<gfxCoreFamily>::appendLaunchKernelWithParams(K
         launchParams.additionalSizeParam,                       // additionalSizeParam
         this->partitionCount,                                   // partitionCount
         launchParams.reserveExtraPayloadSpace,                  // reserveExtraPayloadSpace
+        maxWgCountPerTile,                                      // maxWgCountPerTile
         this->defaultPipelinedThreadArbitrationPolicy,          // defaultPipelinedThreadArbitrationPolicy
         launchParams.isIndirect,                                // isIndirect
         launchParams.isPredicate,                               // isPredicate

level_zero/core/source/kernel/kernel.h

@@ -144,8 +144,7 @@ struct Kernel : _ze_kernel_handle_t, virtual NEO::DispatchKernelEncoderI {
     virtual void patchGlobalOffset() = 0;
     virtual void patchRegionParams(const CmdListKernelLaunchParams &launchParams) = 0;
 
-    virtual ze_result_t suggestMaxCooperativeGroupCount(uint32_t *totalGroupCount, NEO::EngineGroupType engineGroupType,
-                                                        bool isEngineInstanced) = 0;
+    virtual uint32_t suggestMaxCooperativeGroupCount(NEO::EngineGroupType engineGroupType, bool isEngineInstanced, bool forceSingleTileQuery) = 0;
     virtual ze_result_t setCacheConfig(ze_cache_config_flags_t flags) = 0;
 
     virtual ze_result_t getProfileInfo(zet_profile_properties_t *pProfileProperties) = 0;

level_zero/core/source/kernel/kernel_imp.cpp

@@ -477,8 +477,7 @@ ze_result_t KernelImp::suggestGroupSize(uint32_t globalSizeX, uint32_t globalSiz
     return ZE_RESULT_SUCCESS;
 }
 
-ze_result_t KernelImp::suggestMaxCooperativeGroupCount(uint32_t *totalGroupCount, NEO::EngineGroupType engineGroupType,
-                                                       bool isEngineInstanced) {
+uint32_t KernelImp::suggestMaxCooperativeGroupCount(NEO::EngineGroupType engineGroupType, bool isEngineInstanced, bool forceSingleTileQuery) {
     UNRECOVERABLE_IF(0 == groupSize[0]);
     UNRECOVERABLE_IF(0 == groupSize[1]);
     UNRECOVERABLE_IF(0 == groupSize[2]);
@@ -496,20 +495,18 @@ ze_result_t KernelImp::suggestMaxCooperativeGroupCount(uint32_t *totalGroupCount
     bool platformImplicitScaling = helper.platformSupportsImplicitScaling(rootDeviceEnvironment);
     auto deviceBitfield = module->getDevice()->getNEODevice()->getDeviceBitfield();
 
-    if (NEO::ImplicitScalingHelper::isImplicitScalingEnabled(deviceBitfield, platformImplicitScaling)) {
+    if (!forceSingleTileQuery && NEO::ImplicitScalingHelper::isImplicitScalingEnabled(deviceBitfield, platformImplicitScaling)) {
         numSubDevicesForExecution = static_cast<uint32_t>(deviceBitfield.count());
     }
 
-    *totalGroupCount = NEO::KernelHelper::getMaxWorkGroupCount(rootDeviceEnvironment,
-                                                               descriptor,
-                                                               numSubDevicesForExecution,
-                                                               usedSlmSize,
-                                                               workDim,
-                                                               localWorkSize,
-                                                               engineGroupType,
-                                                               isEngineInstanced);
-
-    return ZE_RESULT_SUCCESS;
+    return NEO::KernelHelper::getMaxWorkGroupCount(rootDeviceEnvironment,
+                                                   descriptor,
+                                                   numSubDevicesForExecution,
+                                                   usedSlmSize,
+                                                   workDim,
+                                                   localWorkSize,
+                                                   engineGroupType,
+                                                   isEngineInstanced);
 }
 
 ze_result_t KernelImp::setIndirectAccess(ze_kernel_indirect_access_flags_t flags) {

level_zero/core/source/kernel/kernel_imp.h

@@ -69,8 +69,7 @@ struct KernelImp : Kernel {
 
     ze_result_t getKernelName(size_t *pSize, char *pName) override;
 
-    ze_result_t suggestMaxCooperativeGroupCount(uint32_t *totalGroupCount, NEO::EngineGroupType engineGroupType,
-                                                bool isEngineInstanced) override;
+    uint32_t suggestMaxCooperativeGroupCount(NEO::EngineGroupType engineGroupType, bool isEngineInstanced, bool forceSingleTileQuery) override;
 
     const uint8_t *getCrossThreadData() const override { return crossThreadData.get(); }
     uint32_t getCrossThreadDataSize() const override { return crossThreadDataSize; }

level_zero/core/test/unit_tests/fixtures/kernel_max_cooperative_groups_count_fixture.cpp

@@ -1,5 +1,5 @@
 /*
- * Copyright (C) 2023 Intel Corporation
+ * Copyright (C) 2023-2024 Intel Corporation
  *
  * SPDX-License-Identifier: MIT
  *
@@ -44,8 +44,7 @@ uint32_t KernelImpSuggestMaxCooperativeGroupCountFixture::getMaxWorkGroupCount()
     kernel.groupSize[0] = lws[0];
     kernel.groupSize[1] = lws[1];
     kernel.groupSize[2] = lws[2];
-    uint32_t totalGroupCount = 0;
-    kernel.KernelImp::suggestMaxCooperativeGroupCount(&totalGroupCount, NEO::EngineGroupType::cooperativeCompute, true);
+    uint32_t totalGroupCount = kernel.KernelImp::suggestMaxCooperativeGroupCount(NEO::EngineGroupType::cooperativeCompute, true, false);
     return totalGroupCount;
 }
 } // namespace ult

level_zero/core/test/unit_tests/gen9/test_cmdlist_create_gen9.cpp

@@ -18,6 +18,7 @@
 #include "level_zero/core/test/unit_tests/fixtures/device_fixture.h"
 #include "level_zero/core/test/unit_tests/mocks/mock_cmdlist.h"
 #include "level_zero/core/test/unit_tests/mocks/mock_kernel.h"
+#include "level_zero/core/test/unit_tests/mocks/mock_module.h"
 
 #include <vector>
 
@@ -44,6 +45,7 @@ class CommandListCreateGen9 : public DeviceFixture, public testing::Test {
         DeviceFixture::tearDown();
     }
 
+    std::vector<std::unique_ptr<L0::ult::Module>> mockModules;
     std::vector<void *> isaBuffers;
     ze_group_count_t dispatchKernelArguments;
     void *buffer = nullptr;
@@ -85,6 +87,13 @@ class CommandListCreateGen9 : public DeviceFixture, public testing::Test {
         kernel.perThreadDataSize = perThreadDataSize;
 
         kernel.kernelImmData = &kernelData;
+
+        kernel.groupSize[0] = 1;
+        kernel.groupSize[1] = 1;
+        kernel.groupSize[2] = 1;
+
+        mockModules.emplace_back(std::make_unique<L0::ult::Module>(device, nullptr, ModuleType::builtin));
+        kernel.module = mockModules.back().get();
     }
     void cleanupKernel(WhiteBox<::L0::KernelImmutableData> &kernelData) {
         kernelData.isaGraphicsAllocation.reset(nullptr);

level_zero/core/test/unit_tests/mocks/mock_kernel.cpp

@@ -42,6 +42,10 @@ Mock<::L0::KernelImp>::Mock() : BaseClass() {
     immutableData.kernelDescriptor = &descriptor;
     immutableData.kernelInfo = &info;
     crossThreadData.reset(new uint8_t[100]);
+
+    groupSize[0] = 1;
+    groupSize[1] = 1;
+    groupSize[2] = 1;
 }
 Mock<::L0::KernelImp>::~Mock() {
     delete immutableData.isaGraphicsAllocation.release();