1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
36
37
38
39
40
41
42
43
44
45
46
47
48
49
50
51
52
53
54
55
56
57
58
59
60
61
62
63
64
65
66
67
68
69
70
71
72
73
74
75
76
77
78
79
80
81
82
83
84
85
86
87
88
89
90
91
92
93
94
95
96
97
98
99
100
101
102
103
104
105
106
107
108
109
110
111
112
113
114
115
116
117
118
119
120
121
122
123
124
125
126
127
128
129
130
131
132
133
134
135
136
137
138
139
140
141
|
/*
* Copyright (C) 2021-2025 Intel Corporation
*
* SPDX-License-Identifier: MIT
*
*/
#include "shared/test/common/helpers/debug_manager_state_restore.h"
#include "shared/test/common/helpers/unit_test_helper.h"
#include "shared/test/common/helpers/variable_backup.h"
#include "shared/test/common/mocks/mock_ail_configuration.h"
#include "shared/test/common/mocks/mock_execution_environment.h"
#include "shared/test/common/test_macros/hw_test.h"
namespace NEO {
using IsDG2 = IsProduct<IGFX_DG2>;
using AILTests = ::testing::Test;
TEST(AILTests, whenAILConfigurationCreateFunctionIsCalledWithUnknownGfxCoreThenNullptrIsReturned) {
EXPECT_EQ(nullptr, AILConfiguration::create(IGFX_UNKNOWN));
}
HWTEST2_F(AILTests, givenInitilizedTemplateWhenApplyWithBlenderIsCalledThenFP64SupportIsEnabled, MatchAny) {
AILWhitebox<productFamily> ail;
ail.processName = "blender";
HardwareInfo hwInfo = {};
NEO::RuntimeCapabilityTable &rtTable = hwInfo.capabilityTable;
rtTable.ftrSupportsFP64 = false;
ail.apply(hwInfo);
EXPECT_EQ(rtTable.ftrSupportsFP64, true);
}
HWTEST2_F(AILTests, givenInitilizedTemplateWhenApplyWithAdobePremiereProIsCalledThenPreferredPlatformNameIsSet, MatchAny) {
DebugManagerStateRestore restorer;
AILWhitebox<productFamily> ail;
ail.processName = "Adobe Premiere Pro";
HardwareInfo hwInfo = {};
NEO::RuntimeCapabilityTable &rtTable = hwInfo.capabilityTable;
rtTable.preferredPlatformName = nullptr;
ail.apply(hwInfo);
EXPECT_NE(nullptr, rtTable.preferredPlatformName);
EXPECT_STREQ("Intel(R) OpenCL", rtTable.preferredPlatformName);
}
HWTEST2_F(AILTests, whenCheckingIfSourcesContainKernelThenCorrectResultIsReturned, MatchAny) {
AILWhitebox<productFamily> ail;
std::string kernelSources = R"(
__kernel void CopyBufferToBufferLeftLeftover(
const __global uchar* pSrc,
__global uchar* pDst,
uint srcOffsetInBytes,
uint dstOffsetInBytes)
{
unsigned int gid = get_global_id(0);
pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ];
}
__kernel void CopyBufferToBufferMiddle(
const __global uint* pSrc,
__global uint* pDst,
uint srcOffsetInBytes,
uint dstOffsetInBytes)
{
unsigned int gid = get_global_id(0);
pDst += dstOffsetInBytes >> 2;
pSrc += srcOffsetInBytes >> 2;
uint4 loaded = vload4(gid, pSrc);
vstore4(loaded, gid, pDst);)";
EXPECT_TRUE(ail.sourcesContain(kernelSources, "CopyBufferToBufferMiddle"));
EXPECT_FALSE(ail.sourcesContain(kernelSources, "CopyBufferToBufferMiddleStateless"));
}
HWTEST2_F(AILTests, whenCheckingIsKernelHashCorrectThenCorrectResultIsReturned, MatchAny) {
AILWhitebox<productFamily> ail;
std::string kernelSources = R"(
__kernel void CopyBufferToBufferLeftLeftover(
const __global uchar* pSrc,
__global uchar* pDst,
uint srcOffsetInBytes,
uint dstOffsetInBytes)
{
unsigned int gid = get_global_id(0);
pDst[ gid + dstOffsetInBytes ] = pSrc[ gid + srcOffsetInBytes ];
}
)";
auto expectedHash = 0xafeba928e880fd89;
// If this check fails, probably hash algorithm has been changed.
// In this case we must regenerate hashes in AIL applications kernels
EXPECT_TRUE(ail.isKernelHashCorrect(kernelSources, expectedHash));
kernelSources.insert(0, "text");
EXPECT_FALSE(ail.isKernelHashCorrect(kernelSources, expectedHash));
}
HWTEST2_F(AILTests, whenModifyKernelIfRequiredIsCalledThenDontChangeKernelSources, MatchAny) {
AILWhitebox<productFamily> ail;
std::string kernelSources = "example_kernel(){}";
auto copyKernel = kernelSources;
ail.modifyKernelIfRequired(kernelSources);
EXPECT_STREQ(copyKernel.c_str(), kernelSources.c_str());
}
HWTEST2_F(AILTests, givenAilWhenCheckingContextSyncFlagRequiredThenExpectFalse, MatchAny) {
AILWhitebox<productFamily> ail;
ail.processName = "other";
EXPECT_FALSE(ail.isContextSyncFlagRequired());
}
HWTEST2_F(AILTests, givenAilWhenCheckingOverfetchDisableRequiredThenExpectFalse, MatchAny) {
AILWhitebox<productFamily> ail;
ail.processName = "other";
EXPECT_FALSE(ail.is256BPrefetchDisableRequired());
}
HWTEST2_F(AILTests, givenAilWhenCheckingDrainHostptrsRequiredThenExpectTrue, MatchAny) {
AILWhitebox<productFamily> ail;
ail.processName = "other";
EXPECT_TRUE(ail.drainHostptrs());
}
HWTEST2_F(AILTests, givenAilWhenGetMicrosecondResolutionCalledThenCorrectValueReturned, MatchAny) {
AILWhitebox<productFamily> ail;
EXPECT_EQ(ail.getMicrosecondResolution(), microsecondAdjustment);
}
} // namespace NEO
|
