1 /*
2  * Copyright (C) 2021 Intel Corporation
3  *
4  * SPDX-License-Identifier: MIT
5  *
6  */
7 
8 #include "shared/source/gmm_helper/gmm_helper.h"
9 #include "shared/source/helpers/array_count.h"
10 #include "shared/source/helpers/file_io.h"
11 #include "shared/test/common/helpers/debug_manager_state_restore.h"
12 #include "shared/test/common/helpers/test_files.h"
13 #include "shared/test/common/test_macros/test.h"
14 
15 #include "level_zero/core/source/module/module_imp.h"
16 #include "level_zero/core/test/aub_tests/fixtures/aub_fixture.h"
17 #include "level_zero/core/test/unit_tests/mocks/mock_driver_handle.h"
18 
19 namespace L0 {
20 namespace ult {
21 
22 struct L0BindlessAub : Test<AUBFixtureL0> {
23 
SetUpL0::ult::L0BindlessAub24     void SetUp() override {
25         DebugManager.flags.UseBindlessMode.set(1);
26         DebugManager.flags.UseExternalAllocatorForSshAndDsh.set(1);
27         AUBFixtureL0::SetUp();
28     }
TearDownL0::ult::L0BindlessAub29     void TearDown() override {
30 
31         module->destroy();
32         AUBFixtureL0::TearDown();
33     }
34 
createModuleFromFileL0::ult::L0BindlessAub35     void createModuleFromFile(const std::string &fileName, ze_context_handle_t context, L0::Device *device) {
36         std::string testFile;
37         retrieveBinaryKernelFilenameNoRevision(testFile, fileName + "_", ".bin");
38 
39         size_t size = 0;
40         auto src = loadDataFromFile(
41             testFile.c_str(),
42             size);
43 
44         ASSERT_NE(0u, size);
45         ASSERT_NE(nullptr, src);
46 
47         ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC};
48         moduleDesc.format = ZE_MODULE_FORMAT_NATIVE;
49         moduleDesc.pInputModule = reinterpret_cast<const uint8_t *>(src.get());
50         moduleDesc.inputSize = size;
51         moduleDesc.pBuildFlags = "";
52 
53         module = new ModuleImp(device, nullptr, ModuleType::User);
54         bool success = module->initialize(&moduleDesc, device->getNEODevice());
55         ASSERT_TRUE(success);
56     }
57     DebugManagerStateRestore restorer;
58     ModuleImp *module = nullptr;
59 };
60 
HWTEST2_F(L0BindlessAub,GivenBindlessKernelWhenExecutedThenOutputIsCorrect,IsAtMostXeHpgCore)61 HWTEST2_F(L0BindlessAub, GivenBindlessKernelWhenExecutedThenOutputIsCorrect, IsAtMostXeHpgCore) {
62     constexpr size_t bufferSize = MemoryConstants::pageSize;
63     const uint32_t groupSize[] = {32, 1, 1};
64     const uint32_t groupCount[] = {bufferSize / 32, 1, 1};
65 
66     NEO::DebugManager.flags.UpdateCrossThreadDataSize.set(true);
67 
68     NEO::SVMAllocsManager::UnifiedMemoryProperties unifiedMemoryProperties(InternalMemoryType::HOST_UNIFIED_MEMORY,
69                                                                            context->rootDeviceIndices,
70                                                                            context->deviceBitfields);
71 
72     auto bufferSrc = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(bufferSize, unifiedMemoryProperties);
73     memset(bufferSrc, 55, bufferSize);
74 
75     auto bufferDst = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(bufferSize, unifiedMemoryProperties);
76     memset(bufferDst, 0, bufferSize);
77 
78     auto simulatedCsr = AUBFixtureL0::getSimulatedCsr<FamilyType>();
79     simulatedCsr->initializeEngine();
80 
81     simulatedCsr->writeMemory(*driverHandle->svmAllocsManager->getSVMAlloc(bufferSrc)->gpuAllocations.getDefaultGraphicsAllocation());
82     simulatedCsr->writeMemory(*driverHandle->svmAllocsManager->getSVMAlloc(bufferDst)->gpuAllocations.getDefaultGraphicsAllocation());
83 
84     ze_group_count_t dispatchTraits;
85     dispatchTraits.groupCountX = groupCount[0];
86     dispatchTraits.groupCountY = groupCount[1];
87     dispatchTraits.groupCountZ = groupCount[2];
88 
89     createModuleFromFile("bindless_stateful_copy_buffer", context, device);
90 
91     ze_kernel_handle_t kernel;
92     ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
93     kernelDesc.pKernelName = "StatefulCopyBuffer";
94 
95     EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelCreate(module->toHandle(), &kernelDesc, &kernel));
96     EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetArgumentValue(kernel, 0, sizeof(void *), &bufferSrc));
97     EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetArgumentValue(kernel, 1, sizeof(void *), &bufferDst));
98     EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetGroupSize(kernel, groupSize[0], groupSize[1], groupSize[2]));
99 
100     ze_command_list_handle_t cmdListHandle = commandList->toHandle();
101     EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListAppendLaunchKernel(cmdListHandle, kernel, &dispatchTraits, nullptr, 0, nullptr));
102     commandList->close();
103 
104     pCmdq->executeCommandLists(1, &cmdListHandle, nullptr, false);
105     pCmdq->synchronize(std::numeric_limits<uint32_t>::max());
106 
107     expectMemory<FamilyType>(reinterpret_cast<void *>(driverHandle->svmAllocsManager->getSVMAlloc(bufferDst)->gpuAllocations.getDefaultGraphicsAllocation()->getGpuAddress()),
108                              bufferSrc, bufferSize);
109 
110     EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelDestroy(kernel));
111     driverHandle->svmAllocsManager->freeSVMAlloc(bufferSrc);
112     driverHandle->svmAllocsManager->freeSVMAlloc(bufferDst);
113 }
114 
115 } // namespace ult
116 } // namespace L0