Skip to content

Commit ae340ff

Browse files
JablonskiMateuszCompute-Runtime-Automation
JablonskiMateusz
authored and
Compute-Runtime-Automation
committed
Add L0 aub tests for indirect dispatch
Related-To: NEO-5081 Signed-off-by: Mateusz Jablonski <[email protected]>
1 parent e1ad48c commit ae340ff

File tree

10 files changed

+341
-103
lines changed

10 files changed

+341
-103
lines changed

‎level_zero/core/test/aub_tests/aub_hello_world_test.cpp‎

Lines changed: 2 additions & 11 deletions
Original file line numberDiff line numberDiff line change
@@ -8,23 +8,14 @@
88
#include "test.h"
99

1010
#include "level_zero/core/test/aub_tests/fixtures/aub_fixture.h"
11+
#include "level_zero/core/test/unit_tests/mocks/mock_driver_handle.h"
1112

1213
#include "test_mode.h"
1314

1415
namespace L0 {
1516
namespace ult {
1617

17-
class AUBHelloWorldL0 : public AUBFixtureL0,
18-
public ::testing::Test {
19-
protected:
20-
void SetUp() override {
21-
AUBFixtureL0::SetUp(NEO::defaultHwInfo.get());
22-
}
23-
void TearDown() override {
24-
AUBFixtureL0::TearDown();
25-
}
26-
};
27-
18+
using AUBHelloWorldL0 = Test<AUBFixtureL0>;
2819
TEST_F(AUBHelloWorldL0, whenAppendMemoryCopyIsCalledThenMemoryIsProperlyCopied) {
2920
uint8_t size = 8;
3021
uint8_t val = 255;

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

Lines changed: 10 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,10 @@
1+
#
2+
# Copyright (C) 2021 Intel Corporation
3+
#
4+
# SPDX-License-Identifier: MIT
5+
#
6+
7+
target_sources(ze_intel_gpu_aub_tests PRIVATE
8+
${CMAKE_CURRENT_SOURCE_DIR}/CMakeLists.txt
9+
${CMAKE_CURRENT_SOURCE_DIR}/append_kernel_indirect_aub_tests.cpp
10+
)

‎level_zero/core/test/aub_tests/cmdlist/append_kernel_indirect_aub_tests.cpp‎

Lines changed: 211 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,211 @@
1+
/*
2+
* Copyright (C) 2021 Intel Corporation
3+
*
4+
* SPDX-License-Identifier: MIT
5+
*
6+
*/
7+
8+
#include "shared/source/helpers/array_count.h"
9+
#include "shared/source/helpers/file_io.h"
10+
#include "shared/test/common/helpers/test_files.h"
11+
12+
#include "test.h"
13+
14+
#include "level_zero/core/test/aub_tests/fixtures/aub_fixture.h"
15+
#include "level_zero/core/test/unit_tests/mocks/mock_driver_handle.h"
16+
17+
namespace L0 {
18+
namespace ult {
19+
20+
struct AUBAppendKernelIndirectL0 : Test<AUBFixtureL0> {
21+
22+
static ze_module_handle_t createModuleFromFile(const std::string &fileName, ze_context_handle_t context, ze_device_handle_t device) {
23+
ze_module_handle_t moduleHandle;
24+
std::string testFile;
25+
retrieveBinaryKernelFilenameNoRevision(testFile, fileName + "_", ".bin");
26+
27+
size_t size = 0;
28+
auto src = loadDataFromFile(
29+
testFile.c_str(),
30+
size);
31+
32+
EXPECT_NE(0u, size);
33+
EXPECT_NE(nullptr, src);
34+
35+
if (!src || size == 0) {
36+
return nullptr;
37+
}
38+
39+
ze_module_desc_t moduleDesc = {ZE_STRUCTURE_TYPE_MODULE_DESC};
40+
moduleDesc.format = ZE_MODULE_FORMAT_NATIVE;
41+
moduleDesc.pInputModule = reinterpret_cast<const uint8_t *>(src.get());
42+
moduleDesc.inputSize = size;
43+
moduleDesc.pBuildFlags = "";
44+
45+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeModuleCreate(context, device, &moduleDesc, &moduleHandle, nullptr));
46+
return moduleHandle;
47+
}
48+
};
49+
TEST_F(AUBAppendKernelIndirectL0, whenAppendKernelIndirectThenGlobalWorkSizeIsProperlyProgrammed) {
50+
const uint32_t groupSize[] = {1, 2, 3};
51+
const uint32_t groupCount[] = {4, 3, 1};
52+
const uint32_t expectedGlobalWorkSize[] = {groupSize[0] * groupCount[0],
53+
groupSize[1] * groupCount[1],
54+
groupSize[2] * groupCount[2]};
55+
uint8_t size = 3 * sizeof(uint32_t);
56+
57+
NEO::SVMAllocsManager::UnifiedMemoryProperties unifiedMemoryProperties(InternalMemoryType::HOST_UNIFIED_MEMORY,
58+
context->rootDeviceIndices,
59+
context->deviceBitfields);
60+
61+
auto pDispatchTraits = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(sizeof(ze_group_count_t), unifiedMemoryProperties);
62+
63+
auto outBuffer = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(size, unifiedMemoryProperties);
64+
65+
memset(outBuffer, 0, size);
66+
67+
ze_group_count_t &dispatchTraits = *reinterpret_cast<ze_group_count_t *>(pDispatchTraits);
68+
dispatchTraits.groupCountX = groupCount[0];
69+
dispatchTraits.groupCountY = groupCount[1];
70+
dispatchTraits.groupCountZ = groupCount[2];
71+
72+
ze_module_handle_t moduleHandle = createModuleFromFile("test_kernel", context, device);
73+
ASSERT_NE(nullptr, moduleHandle);
74+
ze_kernel_handle_t kernel;
75+
76+
ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
77+
kernelDesc.pKernelName = "test_get_global_sizes";
78+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelCreate(moduleHandle, &kernelDesc, &kernel));
79+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetArgumentValue(kernel, 0, sizeof(void *), &outBuffer));
80+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetGroupSize(kernel, groupSize[0], groupSize[1], groupSize[2]));
81+
ze_command_list_handle_t cmdListHandle = commandList->toHandle();
82+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListAppendLaunchKernelIndirect(cmdListHandle, kernel, &dispatchTraits, nullptr, 0, nullptr));
83+
commandList->close();
84+
85+
pCmdq->executeCommandLists(1, &cmdListHandle, nullptr, false);
86+
pCmdq->synchronize(std::numeric_limits<uint32_t>::max());
87+
88+
EXPECT_TRUE(csr->expectMemory(outBuffer, expectedGlobalWorkSize, size, AubMemDump::CmdServicesMemTraceMemoryCompare::CompareOperationValues::CompareEqual));
89+
90+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelDestroy(kernel));
91+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeModuleDestroy(moduleHandle));
92+
driverHandle->svmAllocsManager->freeSVMAlloc(outBuffer);
93+
driverHandle->svmAllocsManager->freeSVMAlloc(pDispatchTraits);
94+
}
95+
96+
TEST_F(AUBAppendKernelIndirectL0, whenAppendKernelIndirectThenGroupCountIsProperlyProgrammed) {
97+
const uint32_t groupSize[] = {1, 2, 3};
98+
const uint32_t groupCount[] = {4, 3, 1};
99+
uint8_t size = 3 * sizeof(uint32_t);
100+
101+
NEO::SVMAllocsManager::UnifiedMemoryProperties unifiedMemoryProperties(InternalMemoryType::HOST_UNIFIED_MEMORY,
102+
context->rootDeviceIndices,
103+
context->deviceBitfields);
104+
105+
auto pDispatchTraits = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(sizeof(ze_group_count_t), unifiedMemoryProperties);
106+
107+
auto outBuffer = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(size, unifiedMemoryProperties);
108+
109+
memset(outBuffer, 0, size);
110+
111+
ze_group_count_t &dispatchTraits = *reinterpret_cast<ze_group_count_t *>(pDispatchTraits);
112+
dispatchTraits.groupCountX = groupCount[0];
113+
dispatchTraits.groupCountY = groupCount[1];
114+
dispatchTraits.groupCountZ = groupCount[2];
115+
116+
ze_module_handle_t moduleHandle = createModuleFromFile("test_kernel", context, device);
117+
ASSERT_NE(nullptr, moduleHandle);
118+
ze_kernel_handle_t kernel;
119+
120+
ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
121+
kernelDesc.pKernelName = "test_get_group_count";
122+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelCreate(moduleHandle, &kernelDesc, &kernel));
123+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetArgumentValue(kernel, 0, sizeof(void *), &outBuffer));
124+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetGroupSize(kernel, groupSize[0], groupSize[1], groupSize[2]));
125+
ze_command_list_handle_t cmdListHandle = commandList->toHandle();
126+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListAppendLaunchKernelIndirect(cmdListHandle, kernel, &dispatchTraits, nullptr, 0, nullptr));
127+
commandList->close();
128+
129+
pCmdq->executeCommandLists(1, &cmdListHandle, nullptr, false);
130+
pCmdq->synchronize(std::numeric_limits<uint32_t>::max());
131+
132+
EXPECT_TRUE(csr->expectMemory(outBuffer, groupCount, size, AubMemDump::CmdServicesMemTraceMemoryCompare::CompareOperationValues::CompareEqual));
133+
134+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelDestroy(kernel));
135+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeModuleDestroy(moduleHandle));
136+
driverHandle->svmAllocsManager->freeSVMAlloc(outBuffer);
137+
driverHandle->svmAllocsManager->freeSVMAlloc(pDispatchTraits);
138+
}
139+
140+
TEST_F(AUBAppendKernelIndirectL0, whenAppendKernelIndirectThenWorkDimIsProperlyProgrammed) {
141+
NEO::SVMAllocsManager::UnifiedMemoryProperties unifiedMemoryProperties(InternalMemoryType::HOST_UNIFIED_MEMORY,
142+
context->rootDeviceIndices,
143+
context->deviceBitfields);
144+
145+
std::tuple<std::array<uint32_t, 3> /*groupSize*/, std::array<uint32_t, 3> /*groupCount*/, uint32_t /*expected workdim*/> testData[]{
146+
{{1, 1, 1}, {1, 1, 1}, 1},
147+
{{1, 1, 1}, {2, 1, 1}, 1},
148+
{{1, 1, 1}, {1, 2, 1}, 2},
149+
{{1, 1, 1}, {1, 1, 2}, 3},
150+
{{2, 1, 1}, {1, 1, 1}, 1},
151+
{{2, 1, 1}, {2, 1, 1}, 1},
152+
{{2, 1, 1}, {1, 2, 1}, 2},
153+
{{2, 1, 1}, {1, 1, 2}, 3},
154+
{{1, 2, 1}, {1, 1, 1}, 2},
155+
{{1, 2, 1}, {2, 1, 1}, 2},
156+
{{1, 2, 1}, {1, 2, 1}, 2},
157+
{{1, 2, 1}, {1, 1, 2}, 3},
158+
{{1, 1, 2}, {1, 1, 1}, 3},
159+
{{1, 1, 2}, {2, 1, 1}, 3},
160+
{{1, 1, 2}, {1, 2, 1}, 3},
161+
{{1, 1, 2}, {1, 1, 2}, 3}};
162+
163+
ze_command_list_handle_t cmdListHandle = commandList->toHandle();
164+
ze_module_handle_t moduleHandle = createModuleFromFile("test_kernel", context, device);
165+
ASSERT_NE(nullptr, moduleHandle);
166+
ze_kernel_handle_t kernel;
167+
168+
ze_kernel_desc_t kernelDesc = {ZE_STRUCTURE_TYPE_KERNEL_DESC};
169+
kernelDesc.pKernelName = "test_get_work_dim";
170+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelCreate(moduleHandle, &kernelDesc, &kernel));
171+
for (auto i = 0u; i < arrayCount<>(testData); i++) {
172+
std::array<uint32_t, 3> groupSize;
173+
std::array<uint32_t, 3> groupCount;
174+
uint32_t expectedWorkDim;
175+
176+
std::tie(groupSize, groupCount, expectedWorkDim) = testData[i];
177+
178+
uint8_t size = sizeof(uint32_t);
179+
180+
auto pDispatchTraits = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(sizeof(ze_group_count_t), unifiedMemoryProperties);
181+
182+
auto outBuffer = driverHandle->svmAllocsManager->createHostUnifiedMemoryAllocation(size, unifiedMemoryProperties);
183+
184+
memset(outBuffer, 0, size);
185+
186+
ze_group_count_t &dispatchTraits = *reinterpret_cast<ze_group_count_t *>(pDispatchTraits);
187+
dispatchTraits.groupCountX = groupCount[0];
188+
dispatchTraits.groupCountY = groupCount[1];
189+
dispatchTraits.groupCountZ = groupCount[2];
190+
191+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetArgumentValue(kernel, 0, sizeof(void *), &outBuffer));
192+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelSetGroupSize(kernel, groupSize[0], groupSize[1], groupSize[2]));
193+
194+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeCommandListAppendLaunchKernelIndirect(cmdListHandle, kernel, &dispatchTraits, nullptr, 0, nullptr));
195+
commandList->close();
196+
197+
pCmdq->executeCommandLists(1, &cmdListHandle, nullptr, false);
198+
pCmdq->synchronize(std::numeric_limits<uint32_t>::max());
199+
200+
EXPECT_TRUE(csr->expectMemory(outBuffer, &expectedWorkDim, size, AubMemDump::CmdServicesMemTraceMemoryCompare::CompareOperationValues::CompareEqual));
201+
202+
driverHandle->svmAllocsManager->freeSVMAlloc(outBuffer);
203+
driverHandle->svmAllocsManager->freeSVMAlloc(pDispatchTraits);
204+
commandList->reset();
205+
}
206+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeKernelDestroy(kernel));
207+
EXPECT_EQ(ZE_RESULT_SUCCESS, zeModuleDestroy(moduleHandle));
208+
}
209+
210+
} // namespace ult
211+
} // namespace L0

‎level_zero/core/test/aub_tests/fixtures/aub_fixture.cpp‎

Lines changed: 21 additions & 21 deletions
Original file line numberDiff line numberDiff line change
@@ -5,26 +5,37 @@
55
*
66
*/
77

8-
#include "aub_fixture.h"
8+
#include "level_zero/core/test/aub_tests/fixtures/aub_fixture.h"
99

1010
#include "shared/source/helpers/api_specific_config.h"
11+
#include "shared/test/common/mocks/mock_device.h"
1112

12-
namespace L0 {
13+
#include "level_zero/core/test/unit_tests/mocks/mock_cmdlist.h"
14+
#include "level_zero/core/test/unit_tests/mocks/mock_driver_handle.h"
15+
16+
#include "gtest/gtest.h"
1317

18+
namespace L0 {
19+
AUBFixtureL0::AUBFixtureL0() = default;
20+
AUBFixtureL0::~AUBFixtureL0() = default;
1421
void AUBFixtureL0::prepareCopyEngines(NEO::MockDevice &device, const std::string &filename) {
1522
for (auto i = 0u; i < device.engines.size(); i++) {
16-
if (EngineHelpers::isBcs(device.engines[i].getEngineType())) {
17-
CommandStreamReceiver *pBcsCommandStreamReceiver = nullptr;
18-
pBcsCommandStreamReceiver = AUBCommandStreamReceiver::create(filename, true, *device.executionEnvironment, device.getRootDeviceIndex(), device.getDeviceBitfield());
23+
if (NEO::EngineHelpers::isBcs(device.engines[i].getEngineType())) {
24+
NEO::CommandStreamReceiver *pBcsCommandStreamReceiver = nullptr;
25+
pBcsCommandStreamReceiver = NEO::AUBCommandStreamReceiver::create(filename, true, *device.executionEnvironment, device.getRootDeviceIndex(), device.getDeviceBitfield());
1926
device.resetCommandStreamReceiver(pBcsCommandStreamReceiver, i);
2027
}
2128
}
2229
}
2330

24-
void AUBFixtureL0::SetUp(const HardwareInfo *hardwareInfo) {
25-
const HardwareInfo &hwInfo = hardwareInfo ? *hardwareInfo : *defaultHwInfo;
31+
void AUBFixtureL0::SetUp() {
32+
SetUp(NEO::defaultHwInfo.get());
33+
}
34+
void AUBFixtureL0::SetUp(const NEO::HardwareInfo *hardwareInfo) {
35+
ASSERT_NE(nullptr, hardwareInfo);
36+
const auto &hwInfo = *hardwareInfo;
2637

27-
auto &hwHelper = HwHelper::get(hwInfo.platform.eRenderCoreFamily);
38+
auto &hwHelper = NEO::HwHelper::get(hwInfo.platform.eRenderCoreFamily);
2839
auto engineType = getChosenEngineType(hwInfo);
2940

3041
const ::testing::TestInfo *const testInfo = ::testing::UnitTest::GetInstance()->current_test_info();
@@ -37,9 +48,9 @@ void AUBFixtureL0::SetUp(const HardwareInfo *hardwareInfo) {
3748
executionEnvironment->prepareRootDeviceEnvironments(1u);
3849
executionEnvironment->rootDeviceEnvironments[0]->setHwInfo(&hwInfo);
3950

40-
neoDevice = MockDevice::createWithExecutionEnvironment<MockDevice>(&hwInfo, executionEnvironment, 0u);
51+
neoDevice = NEO::MockDevice::createWithExecutionEnvironment<NEO::MockDevice>(&hwInfo, executionEnvironment, 0u);
4152

42-
this->csr = AUBCommandStreamReceiver::create(strfilename.str(), true, *executionEnvironment, 0, neoDevice->getDeviceBitfield());
53+
this->csr = NEO::AUBCommandStreamReceiver::create(strfilename.str(), true, *executionEnvironment, 0, neoDevice->getDeviceBitfield());
4354
neoDevice->resetCommandStreamReceiver(this->csr);
4455
prepareCopyEngines(*neoDevice, strfilename.str());
4556

@@ -59,17 +70,6 @@ void AUBFixtureL0::SetUp(const HardwareInfo *hardwareInfo) {
5970
ze_result_t returnValue;
6071
commandList.reset(ult::whitebox_cast(CommandList::create(hwInfo.platform.eProductFamily, device, NEO::EngineGroupType::RenderCompute, 0u, returnValue)));
6172

62-
ze_event_pool_desc_t eventPoolDesc = {};
63-
eventPoolDesc.flags = ZE_EVENT_POOL_FLAG_HOST_VISIBLE;
64-
eventPoolDesc.count = 2;
65-
66-
ze_event_desc_t eventDesc = {};
67-
eventDesc.index = 0;
68-
eventDesc.wait = 0;
69-
eventDesc.signal = 0;
70-
71-
eventPool = std::unique_ptr<EventPool>(EventPool::create(device->getDriverHandle(), context, 0, nullptr, &eventPoolDesc));
72-
7373
returnValue = ZE_RESULT_ERROR_UNINITIALIZED;
7474
ze_command_queue_desc_t queueDesc = {};
7575
pCmdq = CommandQueue::create(hwInfo.platform.eProductFamily, device, csr, &queueDesc, false, false, returnValue);

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

Lines changed: 29 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -5,24 +5,37 @@
55
*
66
*/
77

8-
#include "shared/source/helpers/file_io.h"
9-
#include "shared/test/common/helpers/test_files.h"
10-
11-
#include "level_zero/core/source/context/context_imp.h"
12-
#include "level_zero/core/test/unit_tests/fixtures/cmdlist_fixture.h"
13-
#include "level_zero/core/test/unit_tests/mocks/mock_built_ins.h"
14-
#include "level_zero/core/test/unit_tests/mocks/mock_cmdlist.h"
15-
#include "level_zero/core/test/unit_tests/mocks/mock_device.h"
16-
#include "level_zero/core/test/unit_tests/mocks/mock_event.h"
17-
#include "level_zero/core/test/unit_tests/mocks/mock_kernel.h"
18-
19-
#include "third_party/gtest/gtest/gtest.h"
20-
8+
#include <cstdint>
9+
#include <memory>
10+
#include <string>
11+
12+
namespace NEO {
13+
class CommandStreamReceiver;
14+
class MockDevice;
15+
class ExecutionEnvironment;
16+
class MemoryManager;
17+
struct HardwareInfo;
18+
} // namespace NEO
2119
namespace L0 {
20+
namespace ult {
21+
template <typename Type>
22+
struct Mock;
23+
template <typename Type>
24+
struct WhiteBox;
25+
} // namespace ult
26+
27+
struct ContextImp;
28+
struct DriverHandleImp;
29+
struct CommandQueue;
30+
struct CommandList;
31+
struct Device;
2232

2333
class AUBFixtureL0 {
2434
public:
25-
void SetUp(const HardwareInfo *hardwareInfo);
35+
AUBFixtureL0();
36+
virtual ~AUBFixtureL0();
37+
void SetUp();
38+
void SetUp(const NEO::HardwareInfo *hardwareInfo);
2639
void TearDown();
2740
static void prepareCopyEngines(NEO::MockDevice &device, const std::string &filename);
2841

@@ -32,15 +45,13 @@ class AUBFixtureL0 {
3245
NEO::MockDevice *neoDevice = nullptr;
3346

3447
std::unique_ptr<ult::Mock<DriverHandleImp>> driverHandle;
35-
std::unique_ptr<ult::CommandList> commandList;
36-
std::unique_ptr<EventPool> eventPool;
37-
std::unique_ptr<Event> event;
48+
std::unique_ptr<ult::WhiteBox<L0::CommandList>> commandList;
3849

3950
Device *device = nullptr;
4051
ContextImp *context = nullptr;
4152
CommandQueue *pCmdq = nullptr;
4253

43-
CommandStreamReceiver *csr = nullptr;
54+
NEO::CommandStreamReceiver *csr = nullptr;
4455
};
4556

4657
} // namespace L0

0 commit comments

Comments
 (0)