test: multitile sync dispatch aub test

Related-To: NEO-8171

Signed-off-by: Bartosz Dunajski <bartosz.dunajski@intel.com>

level_zero/core/test/aub_tests/cmdlist/CMakeLists.txt

@@ -1,5 +1,5 @@
 #
-# Copyright (C) 2021 Intel Corporation
+# Copyright (C) 2021-2024 Intel Corporation
 #
 # SPDX-License-Identifier: MIT
 #
@@ -7,4 +7,5 @@
 target_sources(ze_intel_gpu_aub_tests PRIVATE
                ${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
                ${CMAKE_CURRENT_SOURCE_DIR}/append_kernel_indirect_aub_tests.cpp
+               ${CMAKE_CURRENT_SOURCE_DIR}/multitile_aub_tests.cpp
 )

level_zero/core/test/aub_tests/cmdlist/multitile_aub_tests.cpp

@@ -0,0 +1,111 @@
+/*
+ * Copyright (C) 2024 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "shared/source/helpers/array_count.h"
+#include "shared/source/helpers/file_io.h"
+#include "shared/test/common/helpers/test_files.h"
+#include "shared/test/common/test_macros/hw_test.h"
+
+#include "level_zero/core/source/cmdqueue/cmdqueue.h"
+#include "level_zero/core/source/context/context_imp.h"
+#include "level_zero/core/source/driver/driver_handle_imp.h"
+#include "level_zero/core/source/module/module.h"
+#include "level_zero/core/test/aub_tests/fixtures/aub_fixture.h"
+#include "level_zero/core/test/aub_tests/fixtures/multicontext_l0_aub_fixture.h"
+#include "level_zero/core/test/unit_tests/mocks/mock_cmdlist.h"
+#include "level_zero/core/test/unit_tests/sources/helper/ze_object_utils.h"
+
+namespace L0 {
+namespace ult {
+
+struct SynchronizedDispatchMultiTileFixture : public MulticontextL0AubFixture {
+    void setUp() {
+        debugManager.flags.ForceSynchronizedDispatchMode.set(1);
+
+        MulticontextL0AubFixture::setUp(2, EnabledCommandStreamers::single, true);
+
+        if (skipped || !rootDevice->isImplicitScalingCapable()) {
+            GTEST_SKIP();
+        }
+
+        ze_context_handle_t hContext;
+        ze_context_desc_t desc = {ZE_STRUCTURE_TYPE_CONTEXT_DESC, nullptr, 0};
+        driverHandle->createContext(&desc, 0u, nullptr, &hContext);
+        ASSERT_NE(nullptr, hContext);
+        context.reset(static_cast<ContextImp *>(Context::fromHandle(hContext)));
+
+        ze_module_handle_t hModule = AUBFixtureL0::createModuleFromFile("test_kernel", context.get(), rootDevice, "");
+        ASSERT_NE(nullptr, hModule);
+        module.reset(Module::fromHandle(hModule));
+
+        ze_kernel_handle_t hKernel;
+        ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
+        kernelDesc.pKernelName = "test_get_global_sizes";
+        zeKernelCreate(hModule, &kernelDesc, &hKernel);
+        ASSERT_NE(nullptr, hKernel);
+        kernel.reset(Kernel::fromHandle(hKernel));
+
+        ze_result_t returnValue;
+        commandList.reset(ult::CommandList::whiteboxCast(CommandList::create(rootDevice->getHwInfo().platform.eProductFamily, rootDevice, NEO::EngineGroupType::compute, 0u, returnValue, false)));
+        ASSERT_NE(nullptr, commandList.get());
+
+        ze_command_queue_desc_t queueDesc = {};
+        cmdQ.reset(CommandQueue::create(rootDevice->getHwInfo().platform.eProductFamily, rootDevice, rootDevice->getNEODevice()->getDefaultEngine().commandStreamReceiver, &queueDesc, false, false, false, returnValue));
+        ASSERT_NE(nullptr, cmdQ.get());
+    }
+
+    DestroyableZeUniquePtr<ContextImp> context;
+    DestroyableZeUniquePtr<Module> module;
+    DestroyableZeUniquePtr<Kernel> kernel;
+    DestroyableZeUniquePtr<ult::WhiteBox<L0::CommandListImp>> commandList;
+    DestroyableZeUniquePtr<CommandQueue> cmdQ;
+};
+
+using SynchronizedDispatchMultiTileL0AubTests = Test<SynchronizedDispatchMultiTileFixture>;
+
+HWTEST_F(SynchronizedDispatchMultiTileL0AubTests, givenFullSyncDispatchWhenExecutingThenDataIsCorrect) {
+    if (!rootDevice->isImplicitScalingCapable()) {
+        GTEST_SKIP();
+    }
+
+    constexpr uint8_t size = 3 * sizeof(uint32_t);
+
+    NEO::SVMAllocsManager::UnifiedMemoryProperties unifiedMemoryProperties(InternalMemoryType::hostUnifiedMemory, 1, context->rootDeviceIndices, context->deviceBitfields);
+
+    auto outBuffer = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(size, unifiedMemoryProperties);
+    memset(outBuffer, 0, size);
+
+    ze_group_count_t groupCount = {};
+    groupCount.groupCountX = 128;
+    groupCount.groupCountY = 1;
+    groupCount.groupCountZ = 1;
+
+    EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetArgumentValue(kernel.get(), 0, sizeof(void *), &outBuffer));
+    EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetGroupSize(kernel.get(), 1, 1, 1));
+    ze_command_list_handle_t cmdListHandle = commandList->toHandle();
+    EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListAppendLaunchKernel(cmdListHandle, kernel.get(), &groupCount, nullptr, 0, nullptr));
+    commandList->close();
+
+    cmdQ->executeCommandLists(1, &cmdListHandle, nullptr, false, nullptr, 0, nullptr);
+    cmdQ->synchronize(std::numeric_limits<uint32_t>::max());
+
+    auto csr = getSimulatedCsr<FamilyType>(0, 0);
+    rootDevice->getNEODevice()->getDefaultEngine().commandStreamReceiver->pollForCompletion();
+
+    const uint32_t expectedGlobalWorkSize[3] = {128, 1, 1};
+    uint64_t expectedTokenValue = 0;
+
+    auto compareEqual = AubMemDump::CmdServicesMemTraceMemoryCompare::CompareOperationValues::CompareEqual;
+
+    EXPECT_TRUE(csr->expectMemory(outBuffer, expectedGlobalWorkSize, size, compareEqual));
+    EXPECT_TRUE(csr->expectMemory(reinterpret_cast<void *>(rootDevice->getSyncDispatchTokenAllocation()->getGpuAddress()), &expectedTokenValue, sizeof(uint64_t), compareEqual));
+
+    driverHandle->svmAllocsManager->freeSVMAlloc(outBuffer);
+}
+
+} // namespace ult
+} // namespace L0

level_zero/core/test/aub_tests/fixtures/CMakeLists.txt

@@ -1,5 +1,5 @@
 #
-# Copyright (C) 2021-2023 Intel Corporation
+# Copyright (C) 2021-2024 Intel Corporation
 #
 # SPDX-License-Identifier: MIT
 #
@@ -9,6 +9,8 @@ target_sources(ze_intel_gpu_aub_tests PRIVATE
                ${CMAKE_CURRENT_SOURCE_DIR}/aub_fixture.cpp
                ${CMAKE_CURRENT_SOURCE_DIR}/aub_fixture.h
                ${CMAKE_CURRENT_SOURCE_DIR}/create_aub_ult_builtin_functions_lib.cpp
+               ${CMAKE_CURRENT_SOURCE_DIR}/multicontext_l0_aub_fixture.h
+               ${CMAKE_CURRENT_SOURCE_DIR}/multicontext_l0_aub_fixture.cpp
 
 )
 add_subdirectories()

level_zero/core/test/aub_tests/fixtures/aub_fixture.h

@@ -101,7 +101,7 @@ class AUBFixtureL0 {
         }
     }
 
-    ze_module_handle_t createModuleFromFile(const std::string &fileName, ze_context_handle_t context, ze_device_handle_t device, const std::string &buildFlags);
+    static ze_module_handle_t createModuleFromFile(const std::string &fileName, ze_context_handle_t context, ze_device_handle_t device, const std::string &buildFlags);
 
     std::string aubFileName;
     std::unique_ptr<VariableBackup<NEO::UltHwConfig>> backupUltConfig;

level_zero/core/test/aub_tests/fixtures/multicontext_l0_aub_fixture.cpp

@@ -0,0 +1,62 @@
+/*
+ * Copyright (C) 2024 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#include "level_zero/core/test/aub_tests/fixtures/multicontext_l0_aub_fixture.h"
+
+#include "shared/source/helpers/gfx_core_helper.h"
+#include "shared/source/helpers/hw_info.h"
+#include "shared/test/common/helpers/ult_hw_config.h"
+#include "shared/test/common/helpers/variable_backup.h"
+#include "shared/test/common/mocks/mock_device.h"
+#include "shared/test/common/test_macros/test.h"
+
+#include "level_zero/core/source/device/device_imp.h"
+
+void MulticontextL0AubFixture::setUp(uint32_t numberOfTiles, EnabledCommandStreamers enabledCommandStreamers, bool implicitScaling) {
+    if (implicitScaling) {
+        debugManager.flags.EnableImplicitScaling.set(1);
+    }
+    MulticontextAubFixture::setUp(numberOfTiles, enabledCommandStreamers, false);
+
+    if (skipped) {
+        GTEST_SKIP();
+    }
+}
+
+CommandStreamReceiver *MulticontextL0AubFixture::getGpgpuCsr(uint32_t tile, uint32_t engine) {
+    return subDevices[tile]->getNEODevice()->getEngine(engine).commandStreamReceiver;
+}
+
+void MulticontextL0AubFixture::createDevices(const HardwareInfo &hwInfo, uint32_t numTiles) {
+    VariableBackup<UltHwConfig> backup(&ultHwConfig);
+    ultHwConfig.useHwCsr = true;
+
+    auto executionEnvironment = new NEO::ExecutionEnvironment();
+    executionEnvironment->prepareRootDeviceEnvironments(1u);
+    executionEnvironment->rootDeviceEnvironments[0]->setHwInfoAndInitHelpers(&hwInfo);
+    executionEnvironment->rootDeviceEnvironments[0]->initGmm();
+    executionEnvironment->calculateMaxOsContextCount();
+
+    auto neoDevice = NEO::MockDevice::createWithExecutionEnvironment<NEO::MockDevice>(&hwInfo, executionEnvironment, 0u);
+
+    NEO::DeviceVector devices;
+    devices.push_back(std::unique_ptr<NEO::Device>(neoDevice));
+    driverHandle = std::make_unique<L0::ult::Mock<L0::DriverHandleImp>>();
+
+    driverHandle->initialize(std::move(devices));
+
+    rootDevice = driverHandle->devices[0];
+
+    std::vector<ze_device_handle_t> subDevicesH;
+    subDevicesH.resize(numTiles);
+
+    EXPECT_EQ(ZE_RESULT_SUCCESS, rootDevice->getSubDevices(&numberOfEnabledTiles, subDevicesH.data()));
+
+    for (uint32_t i = 0; i < numTiles; i++) {
+        subDevices.push_back(L0::Device::fromHandle(subDevicesH[i]));
+    }
+}

level_zero/core/test/aub_tests/fixtures/multicontext_l0_aub_fixture.h

@@ -0,0 +1,41 @@
+/*
+ * Copyright (C) 2024 Intel Corporation
+ *
+ * SPDX-License-Identifier: MIT
+ *
+ */
+
+#pragma once
+
+#include "shared/test/common/fixtures/aub_fixtures/multicontext_aub_fixture.h"
+
+#include "level_zero/core/test/unit_tests/mock.h"
+#include "level_zero/core/test/unit_tests/mocks/mock_driver_handle.h"
+
+#include <memory>
+#include <vector>
+
+namespace L0 {
+struct Device;
+struct DriverHandleImp;
+;
+
+namespace ult {
+template <typename Type>
+struct Mock;
+template <typename Type>
+struct WhiteBox;
+} // namespace ult
+} // namespace L0
+
+struct MulticontextL0AubFixture : public NEO::MulticontextAubFixture {
+    void setUp(uint32_t numberOfTiles, EnabledCommandStreamers enabledCommandStreamers, bool implicitScaling);
+
+    CommandStreamReceiver *getGpgpuCsr(uint32_t tile, uint32_t engine) override;
+    void createDevices(const NEO::HardwareInfo &hwInfo, uint32_t numTiles) override;
+
+    std::unique_ptr<L0::ult::Mock<L0::DriverHandleImp>> driverHandle;
+    NEO::ExecutionEnvironment *executionEnvironment = nullptr;
+    std::vector<L0::Device *> subDevices;
+    L0::Device *rootDevice = nullptr;
+};