1 /*
2 * Copyright (C) 2018-2021 Intel Corporation
3 *
4 * SPDX-License-Identifier: MIT
5 *
6 */
7
8 #include "opencl/test/unit_test/program/program_tests.h"
9
10 #include "shared/source/command_stream/command_stream_receiver_hw.h"
11 #include "shared/source/compiler_interface/compiler_warnings/compiler_warnings.h"
12 #include "shared/source/compiler_interface/intermediate_representations.h"
13 #include "shared/source/device_binary_format/elf/elf_decoder.h"
14 #include "shared/source/device_binary_format/elf/ocl_elf.h"
15 #include "shared/source/device_binary_format/patchtokens_decoder.h"
16 #include "shared/source/gmm_helper/gmm_helper.h"
17 #include "shared/source/helpers/aligned_memory.h"
18 #include "shared/source/helpers/hash.h"
19 #include "shared/source/helpers/hw_helper.h"
20 #include "shared/source/helpers/ptr_math.h"
21 #include "shared/source/helpers/string.h"
22 #include "shared/source/memory_manager/allocations_list.h"
23 #include "shared/source/memory_manager/graphics_allocation.h"
24 #include "shared/source/memory_manager/surface.h"
25 #include "shared/source/os_interface/os_context.h"
26 #include "shared/test/common/helpers/debug_manager_state_restore.h"
27 #include "shared/test/common/helpers/kernel_binary_helper.h"
28 #include "shared/test/common/libult/global_environment.h"
29 #include "shared/test/common/libult/ult_command_stream_receiver.h"
30 #include "shared/test/common/mocks/mock_allocation_properties.h"
31 #include "shared/test/common/mocks/mock_compiler_interface.h"
32 #include "shared/test/common/mocks/mock_graphics_allocation.h"
33 #include "shared/test/common/test_macros/test.h"
34 #include "shared/test/unit_test/device_binary_format/patchtokens_tests.h"
35 #include "shared/test/unit_test/device_binary_format/zebin_tests.h"
36 #include "shared/test/unit_test/utilities/base_object_utils.h"
37
38 #include "opencl/source/gtpin/gtpin_notify.h"
39 #include "opencl/source/helpers/hardware_commands_helper.h"
40 #include "opencl/source/kernel/kernel.h"
41 #include "opencl/source/program/create.inl"
42 #include "opencl/test/unit_test/fixtures/cl_device_fixture.h"
43 #include "opencl/test/unit_test/fixtures/multi_root_device_fixture.h"
44 #include "opencl/test/unit_test/mocks/mock_kernel.h"
45 #include "opencl/test/unit_test/mocks/mock_platform.h"
46 #include "opencl/test/unit_test/mocks/mock_program.h"
47 #include "opencl/test/unit_test/program/program_from_binary.h"
48 #include "opencl/test/unit_test/program/program_with_source.h"
49 #include "opencl/test/unit_test/test_macros/test_checks_ocl.h"
50
51 #include "compiler_options.h"
52 #include "gmock/gmock.h"
53 #include "gtest/gtest.h"
54
55 #include <map>
56 #include <memory>
57 #include <string>
58 #include <vector>
59
60 using namespace NEO;
61
SetUp()62 void ProgramTests::SetUp() {
63 ClDeviceFixture::SetUp();
64 cl_device_id device = pClDevice;
65 ContextFixture::SetUp(1, &device);
66 }
67
TearDown()68 void ProgramTests::TearDown() {
69 ContextFixture::TearDown();
70 ClDeviceFixture::TearDown();
71 }
72
73 class NoCompilerInterfaceRootDeviceEnvironment : public RootDeviceEnvironment {
74 public:
NoCompilerInterfaceRootDeviceEnvironment(ExecutionEnvironment & executionEnvironment)75 NoCompilerInterfaceRootDeviceEnvironment(ExecutionEnvironment &executionEnvironment) : RootDeviceEnvironment(executionEnvironment) {
76 *hwInfo = *defaultHwInfo;
77 }
78
getCompilerInterface()79 CompilerInterface *getCompilerInterface() override {
80 return nullptr;
81 }
82
initAilConfiguration()83 bool initAilConfiguration() override {
84 return true;
85 }
86 };
87
88 class FailingGenBinaryProgram : public MockProgram {
89 public:
90 using MockProgram::MockProgram;
processGenBinary(const ClDevice & clDevice)91 cl_int processGenBinary(const ClDevice &clDevice) override { return CL_INVALID_BINARY; }
92 };
93
94 class SucceedingGenBinaryProgram : public MockProgram {
95 public:
96 using MockProgram::MockProgram;
processGenBinary(const ClDevice & clDevice)97 cl_int processGenBinary(const ClDevice &clDevice) override { return CL_SUCCESS; }
98 };
99
100 using ProgramFromBinaryTest = ProgramFromBinaryFixture;
101
TEST_F(ProgramFromBinaryTest,WhenBuildingProgramThenSuccessIsReturned)102 TEST_F(ProgramFromBinaryTest, WhenBuildingProgramThenSuccessIsReturned) {
103 retVal = pProgram->build(
104 pProgram->getDevices(),
105 nullptr,
106 false);
107
108 EXPECT_EQ(CL_SUCCESS, retVal);
109 }
110
TEST_F(ProgramFromBinaryTest,WhenGettingProgramContextInfoThenCorrectContextIsReturned)111 TEST_F(ProgramFromBinaryTest, WhenGettingProgramContextInfoThenCorrectContextIsReturned) {
112 cl_context contextRet = reinterpret_cast<cl_context>(static_cast<uintptr_t>(0xdeaddead));
113 size_t paramValueSizeRet = 0;
114
115 retVal = pProgram->getInfo(
116 CL_PROGRAM_CONTEXT,
117 sizeof(cl_context),
118 &contextRet,
119 ¶mValueSizeRet);
120
121 EXPECT_EQ(CL_SUCCESS, retVal);
122 EXPECT_EQ(pContext, contextRet);
123 EXPECT_EQ(sizeof(cl_context), paramValueSizeRet);
124 }
125
TEST_F(ProgramFromBinaryTest,GivenNonNullParamValueWhenGettingProgramBinaryInfoThenCorrectBinaryIsReturned)126 TEST_F(ProgramFromBinaryTest, GivenNonNullParamValueWhenGettingProgramBinaryInfoThenCorrectBinaryIsReturned) {
127 size_t paramValueSize = sizeof(unsigned char **);
128 size_t paramValueSizeRet = 0;
129 auto testBinary = std::make_unique<char[]>(knownSourceSize);
130
131 retVal = pProgram->getInfo(
132 CL_PROGRAM_BINARIES,
133 paramValueSize,
134 &testBinary,
135 ¶mValueSizeRet);
136
137 EXPECT_EQ(CL_SUCCESS, retVal);
138 EXPECT_EQ(paramValueSize, paramValueSizeRet);
139 EXPECT_STREQ((const char *)knownSource.get(), (const char *)testBinary.get());
140 }
141
TEST_F(ProgramFromBinaryTest,GivenNullParamValueWhenGettingProgramBinaryInfoThenSuccessIsReturned)142 TEST_F(ProgramFromBinaryTest, GivenNullParamValueWhenGettingProgramBinaryInfoThenSuccessIsReturned) {
143 size_t paramValueSize = sizeof(unsigned char **);
144 size_t paramValueSizeRet = 0;
145
146 retVal = pProgram->getInfo(
147 CL_PROGRAM_BINARIES,
148 0,
149 nullptr,
150 ¶mValueSizeRet);
151 EXPECT_EQ(CL_SUCCESS, retVal);
152 EXPECT_EQ(paramValueSize, paramValueSizeRet);
153 }
154
TEST_F(ProgramFromBinaryTest,GivenNonNullParamValueAndParamValueSizeZeroWhenGettingProgramBinaryInfoThenInvalidValueErrorIsReturned)155 TEST_F(ProgramFromBinaryTest, GivenNonNullParamValueAndParamValueSizeZeroWhenGettingProgramBinaryInfoThenInvalidValueErrorIsReturned) {
156 size_t paramValueSizeRet = 0;
157 auto testBinary = std::make_unique<char[]>(knownSourceSize);
158
159 retVal = pProgram->getInfo(
160 CL_PROGRAM_BINARIES,
161 0,
162 &testBinary,
163 ¶mValueSizeRet);
164 EXPECT_EQ(CL_INVALID_VALUE, retVal);
165 }
166
TEST_F(ProgramFromBinaryTest,GivenInvalidParametersWhenGettingProgramInfoThenValueSizeRetIsNotUpdated)167 TEST_F(ProgramFromBinaryTest, GivenInvalidParametersWhenGettingProgramInfoThenValueSizeRetIsNotUpdated) {
168 size_t paramValueSizeRet = 0x1234;
169 auto testBinary = std::make_unique<char[]>(knownSourceSize);
170
171 retVal = pProgram->getInfo(
172 CL_PROGRAM_BINARIES,
173 0,
174 &testBinary,
175 ¶mValueSizeRet);
176 EXPECT_EQ(CL_INVALID_VALUE, retVal);
177 EXPECT_EQ(0x1234u, paramValueSizeRet);
178 }
179
TEST_F(ProgramFromBinaryTest,GivenInvalidParamWhenGettingProgramBinaryInfoThenInvalidValueErrorIsReturned)180 TEST_F(ProgramFromBinaryTest, GivenInvalidParamWhenGettingProgramBinaryInfoThenInvalidValueErrorIsReturned) {
181 size_t paramValueSizeRet = 0;
182 auto testBinary = std::make_unique<char[]>(knownSourceSize);
183
184 retVal = pProgram->getInfo(
185 CL_PROGRAM_BUILD_STATUS,
186 0,
187 nullptr,
188 ¶mValueSizeRet);
189 EXPECT_EQ(CL_INVALID_VALUE, retVal);
190 }
191
TEST_F(ProgramFromBinaryTest,WhenGettingBinarySizesThenCorrectSizesAreReturned)192 TEST_F(ProgramFromBinaryTest, WhenGettingBinarySizesThenCorrectSizesAreReturned) {
193 size_t paramValueSize = sizeof(size_t *);
194 size_t paramValue[1];
195 size_t paramValueSizeRet = 0;
196
197 retVal = pProgram->getInfo(
198 CL_PROGRAM_BINARY_SIZES,
199 paramValueSize,
200 paramValue,
201 ¶mValueSizeRet);
202
203 EXPECT_EQ(CL_SUCCESS, retVal);
204 EXPECT_EQ(knownSourceSize, paramValue[0]);
205 EXPECT_EQ(paramValueSize, paramValueSizeRet);
206 }
207
TEST_F(ProgramFromBinaryTest,GivenProgramWithOneKernelWhenGettingNumKernelsThenOneIsReturned)208 TEST_F(ProgramFromBinaryTest, GivenProgramWithOneKernelWhenGettingNumKernelsThenOneIsReturned) {
209 size_t paramValue = 0;
210 size_t paramValueSize = sizeof(paramValue);
211 size_t paramValueSizeRet = 0;
212
213 retVal = pProgram->build(
214 pProgram->getDevices(),
215 nullptr,
216 false);
217 ASSERT_EQ(CL_SUCCESS, retVal);
218
219 retVal = pProgram->getInfo(
220 CL_PROGRAM_NUM_KERNELS,
221 paramValueSize,
222 ¶mValue,
223 ¶mValueSizeRet);
224
225 EXPECT_EQ(CL_SUCCESS, retVal);
226 EXPECT_EQ(1u, paramValue);
227 EXPECT_EQ(paramValueSize, paramValueSizeRet);
228 }
229
TEST_F(ProgramFromBinaryTest,GivenProgramWithNoExecutableCodeWhenGettingNumKernelsThenInvalidProgramExecutableErrorIsReturned)230 TEST_F(ProgramFromBinaryTest, GivenProgramWithNoExecutableCodeWhenGettingNumKernelsThenInvalidProgramExecutableErrorIsReturned) {
231 size_t paramValue = 0;
232 size_t paramValueSize = sizeof(paramValue);
233 size_t paramValueSizeRet = 0;
234
235 CreateProgramFromBinary(pContext, pContext->getDevices(), binaryFileName);
236 MockProgram *p = pProgram;
237 p->setBuildStatus(CL_BUILD_NONE);
238
239 retVal = pProgram->getInfo(
240 CL_PROGRAM_NUM_KERNELS,
241 paramValueSize,
242 ¶mValue,
243 ¶mValueSizeRet);
244 EXPECT_EQ(CL_INVALID_PROGRAM_EXECUTABLE, retVal);
245 }
246
TEST_F(ProgramFromBinaryTest,WhenGettingKernelNamesThenCorrectNameIsReturned)247 TEST_F(ProgramFromBinaryTest, WhenGettingKernelNamesThenCorrectNameIsReturned) {
248 size_t paramValueSize = sizeof(size_t *);
249 size_t paramValueSizeRet = 0;
250
251 retVal = pProgram->build(
252 pProgram->getDevices(),
253 nullptr,
254 false);
255 ASSERT_EQ(CL_SUCCESS, retVal);
256
257 // get info successfully about required sizes for kernel names
258 retVal = pProgram->getInfo(
259 CL_PROGRAM_KERNEL_NAMES,
260 0,
261 nullptr,
262 ¶mValueSizeRet);
263 ASSERT_EQ(CL_SUCCESS, retVal);
264 ASSERT_NE(0u, paramValueSizeRet);
265
266 // get info successfully about kernel names
267 auto paramValue = std::make_unique<char[]>(paramValueSizeRet);
268 paramValueSize = paramValueSizeRet;
269 ASSERT_NE(paramValue, nullptr);
270
271 size_t expectedKernelsStringSize = strlen(kernelName) + 1;
272 retVal = pProgram->getInfo(
273 CL_PROGRAM_KERNEL_NAMES,
274 paramValueSize,
275 paramValue.get(),
276 ¶mValueSizeRet);
277
278 EXPECT_EQ(CL_SUCCESS, retVal);
279 EXPECT_STREQ(kernelName, (char *)paramValue.get());
280 EXPECT_EQ(expectedKernelsStringSize, paramValueSizeRet);
281 }
282
TEST_F(ProgramFromBinaryTest,GivenProgramWithNoExecutableCodeWhenGettingKernelNamesThenInvalidProgramExecutableErrorIsReturned)283 TEST_F(ProgramFromBinaryTest, GivenProgramWithNoExecutableCodeWhenGettingKernelNamesThenInvalidProgramExecutableErrorIsReturned) {
284 size_t paramValueSize = sizeof(size_t *);
285 size_t paramValueSizeRet = 0;
286
287 CreateProgramFromBinary(pContext, pContext->getDevices(), binaryFileName);
288 MockProgram *p = pProgram;
289 p->setBuildStatus(CL_BUILD_NONE);
290
291 retVal = pProgram->getInfo(
292 CL_PROGRAM_KERNEL_NAMES,
293 paramValueSize,
294 nullptr,
295 ¶mValueSizeRet);
296 EXPECT_EQ(CL_INVALID_PROGRAM_EXECUTABLE, retVal);
297 }
298
TEST_F(ProgramFromBinaryTest,WhenGettingProgramScopeGlobalCtorsAndDtorsPresentInfoThenCorrectValueIsReturned)299 TEST_F(ProgramFromBinaryTest, WhenGettingProgramScopeGlobalCtorsAndDtorsPresentInfoThenCorrectValueIsReturned) {
300 cl_uint paramRet = 0;
301 cl_uint expectedParam = CL_FALSE;
302 size_t paramSizeRet = 0;
303
304 retVal = pProgram->getInfo(
305 CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT,
306 sizeof(cl_uint),
307 ¶mRet,
308 ¶mSizeRet);
309
310 EXPECT_EQ(CL_SUCCESS, retVal);
311 EXPECT_EQ(sizeof(cl_uint), paramSizeRet);
312 EXPECT_EQ(expectedParam, paramRet);
313
314 retVal = pProgram->getInfo(
315 CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT,
316 sizeof(cl_uint),
317 ¶mRet,
318 ¶mSizeRet);
319
320 EXPECT_EQ(CL_SUCCESS, retVal);
321 EXPECT_EQ(sizeof(cl_uint), paramSizeRet);
322 EXPECT_EQ(expectedParam, paramRet);
323 }
324
TEST_F(ProgramFromBinaryTest,GivenNullDeviceWhenGettingBuildStatusThenBuildNoneIsReturned)325 TEST_F(ProgramFromBinaryTest, GivenNullDeviceWhenGettingBuildStatusThenBuildNoneIsReturned) {
326 cl_device_id device = pClDevice;
327 cl_build_status buildStatus = 0;
328 size_t paramValueSize = sizeof(buildStatus);
329 size_t paramValueSizeRet = 0;
330
331 retVal = pProgram->getBuildInfo(
332 device,
333 CL_PROGRAM_BUILD_STATUS,
334 paramValueSize,
335 &buildStatus,
336 ¶mValueSizeRet);
337
338 EXPECT_EQ(CL_SUCCESS, retVal);
339 EXPECT_EQ(paramValueSize, paramValueSizeRet);
340 EXPECT_EQ(CL_BUILD_NONE, buildStatus);
341 }
342
TEST_F(ProgramFromBinaryTest,GivenInvalidParametersWhenGettingBuildInfoThenValueSizeRetIsNotUpdated)343 TEST_F(ProgramFromBinaryTest, GivenInvalidParametersWhenGettingBuildInfoThenValueSizeRetIsNotUpdated) {
344 cl_device_id device = pClDevice;
345 cl_build_status buildStatus = 0;
346 size_t paramValueSize = sizeof(buildStatus);
347 size_t paramValueSizeRet = 0x1234;
348
349 retVal = pProgram->getBuildInfo(
350 device,
351 0,
352 paramValueSize,
353 &buildStatus,
354 ¶mValueSizeRet);
355
356 EXPECT_EQ(CL_INVALID_VALUE, retVal);
357 EXPECT_EQ(0x1234u, paramValueSizeRet);
358 }
359
TEST_F(ProgramFromBinaryTest,GivenDefaultDeviceWhenGettingBuildOptionsThenBuildOptionsAreEmpty)360 TEST_F(ProgramFromBinaryTest, GivenDefaultDeviceWhenGettingBuildOptionsThenBuildOptionsAreEmpty) {
361 cl_device_id device = pClDevice;
362 size_t paramValueSizeRet = 0u;
363 size_t paramValueSize = 0u;
364
365 retVal = pProgram->getBuildInfo(
366 device,
367 CL_PROGRAM_BUILD_OPTIONS,
368 0,
369 nullptr,
370 ¶mValueSizeRet);
371
372 EXPECT_EQ(CL_SUCCESS, retVal);
373 EXPECT_NE(paramValueSizeRet, 0u);
374
375 auto paramValue = std::make_unique<char[]>(paramValueSizeRet);
376 paramValueSize = paramValueSizeRet;
377
378 retVal = pProgram->getBuildInfo(
379 device,
380 CL_PROGRAM_BUILD_OPTIONS,
381 paramValueSize,
382 paramValue.get(),
383 ¶mValueSizeRet);
384
385 EXPECT_EQ(CL_SUCCESS, retVal);
386 EXPECT_STREQ("", (char *)paramValue.get());
387 }
388
TEST_F(ProgramFromBinaryTest,GivenDefaultDeviceWhenGettingLogThenLogEmpty)389 TEST_F(ProgramFromBinaryTest, GivenDefaultDeviceWhenGettingLogThenLogEmpty) {
390 cl_device_id device = pClDevice;
391 size_t paramValueSizeRet = 0u;
392 size_t paramValueSize = 0u;
393
394 retVal = pProgram->getBuildInfo(
395 device,
396 CL_PROGRAM_BUILD_LOG,
397 0,
398 nullptr,
399 ¶mValueSizeRet);
400
401 EXPECT_EQ(CL_SUCCESS, retVal);
402 EXPECT_NE(paramValueSizeRet, 0u);
403
404 auto paramValue = std::make_unique<char[]>(paramValueSizeRet);
405 paramValueSize = paramValueSizeRet;
406
407 retVal = pProgram->getBuildInfo(
408 device,
409 CL_PROGRAM_BUILD_LOG,
410 paramValueSize,
411 paramValue.get(),
412 ¶mValueSizeRet);
413
414 EXPECT_EQ(CL_SUCCESS, retVal);
415 EXPECT_STREQ("", (char *)paramValue.get());
416 }
417
TEST_F(ProgramFromBinaryTest,GivenLogEntriesWhenGetBuildLogThenLogIsApended)418 TEST_F(ProgramFromBinaryTest, GivenLogEntriesWhenGetBuildLogThenLogIsApended) {
419 cl_device_id device = pClDevice;
420 size_t paramValueSizeRet = 0u;
421 size_t paramValueSize = 0u;
422
423 retVal = pProgram->getBuildInfo(
424 device,
425 CL_PROGRAM_BUILD_LOG,
426 0,
427 nullptr,
428 ¶mValueSizeRet);
429
430 EXPECT_EQ(CL_SUCCESS, retVal);
431 EXPECT_NE(paramValueSizeRet, 0u);
432
433 auto paramValue = std::make_unique<char[]>(paramValueSizeRet);
434 paramValueSize = paramValueSizeRet;
435
436 retVal = pProgram->getBuildInfo(
437 device,
438 CL_PROGRAM_BUILD_LOG,
439 paramValueSize,
440 paramValue.get(),
441 ¶mValueSizeRet);
442
443 EXPECT_EQ(CL_SUCCESS, retVal);
444 EXPECT_STREQ("", (char *)paramValue.get());
445
446 // Add more text to the log
447 pProgram->updateBuildLog(pClDevice->getRootDeviceIndex(), "testing", 8);
448 pProgram->updateBuildLog(pClDevice->getRootDeviceIndex(), "several", 8);
449
450 retVal = pProgram->getBuildInfo(
451 device,
452 CL_PROGRAM_BUILD_LOG,
453 0,
454 nullptr,
455 ¶mValueSizeRet);
456
457 EXPECT_EQ(CL_SUCCESS, retVal);
458 EXPECT_GE(paramValueSizeRet, 16u);
459 paramValue = std::make_unique<char[]>(paramValueSizeRet);
460
461 paramValueSize = paramValueSizeRet;
462
463 retVal = pProgram->getBuildInfo(
464 device,
465 CL_PROGRAM_BUILD_LOG,
466 paramValueSize,
467 paramValue.get(),
468 ¶mValueSizeRet);
469
470 EXPECT_EQ(CL_SUCCESS, retVal);
471
472 EXPECT_NE(nullptr, strstr(paramValue.get(), "testing"));
473
474 const char *paramValueContinued = strstr(paramValue.get(), "testing") + 7;
475 ASSERT_NE(nullptr, strstr(paramValueContinued, "several"));
476 }
477
TEST_F(ProgramFromBinaryTest,GivenNullParamValueWhenGettingProgramBinaryTypeThenParamValueSizeIsReturned)478 TEST_F(ProgramFromBinaryTest, GivenNullParamValueWhenGettingProgramBinaryTypeThenParamValueSizeIsReturned) {
479 cl_device_id device = pClDevice;
480 size_t paramValueSizeRet = 0u;
481 size_t paramValueSize = 0u;
482
483 retVal = pProgram->getBuildInfo(
484 device,
485 CL_PROGRAM_BINARY_TYPE,
486 paramValueSize,
487 nullptr,
488 ¶mValueSizeRet);
489
490 EXPECT_EQ(CL_SUCCESS, retVal);
491 EXPECT_NE(paramValueSizeRet, 0u);
492 }
493
TEST_F(ProgramFromBinaryTest,WhenGettingProgramBinaryTypeThenCorrectProgramTypeIsReturned)494 TEST_F(ProgramFromBinaryTest, WhenGettingProgramBinaryTypeThenCorrectProgramTypeIsReturned) {
495 cl_device_id device = pClDevice;
496 cl_program_binary_type programType = 0;
497 char *paramValue = (char *)&programType;
498 size_t paramValueSizeRet = 0u;
499 size_t paramValueSize = 0u;
500
501 retVal = pProgram->getBuildInfo(
502 device,
503 CL_PROGRAM_BINARY_TYPE,
504 paramValueSize,
505 nullptr,
506 ¶mValueSizeRet);
507
508 EXPECT_EQ(CL_SUCCESS, retVal);
509 EXPECT_NE(paramValueSizeRet, 0u);
510
511 paramValueSize = paramValueSizeRet;
512 retVal = pProgram->getBuildInfo(
513 device,
514 CL_PROGRAM_BINARY_TYPE,
515 paramValueSize,
516 paramValue,
517 ¶mValueSizeRet);
518 EXPECT_EQ(CL_SUCCESS, retVal);
519 EXPECT_EQ((cl_program_binary_type)CL_PROGRAM_BINARY_TYPE_EXECUTABLE, programType);
520 }
521
TEST_F(ProgramFromBinaryTest,GivenInvalidParamWhenGettingBuildInfoThenInvalidValueErrorIsReturned)522 TEST_F(ProgramFromBinaryTest, GivenInvalidParamWhenGettingBuildInfoThenInvalidValueErrorIsReturned) {
523 cl_device_id device = pClDevice;
524 size_t paramValueSizeRet = 0u;
525
526 retVal = pProgram->getBuildInfo(
527 device,
528 CL_PROGRAM_KERNEL_NAMES,
529 0,
530 nullptr,
531 ¶mValueSizeRet);
532 EXPECT_EQ(CL_INVALID_VALUE, retVal);
533 }
534
TEST_F(ProgramFromBinaryTest,GivenGlobalVariableTotalSizeSetWhenGettingBuildGlobalVariableTotalSizeThenCorrectSizeIsReturned)535 TEST_F(ProgramFromBinaryTest, GivenGlobalVariableTotalSizeSetWhenGettingBuildGlobalVariableTotalSizeThenCorrectSizeIsReturned) {
536 cl_device_id device = pClDevice;
537 size_t globalVarSize = 22;
538 size_t paramValueSize = sizeof(globalVarSize);
539 size_t paramValueSizeRet = 0;
540 char *paramValue = (char *)&globalVarSize;
541
542 // get build info as is
543 retVal = pProgram->getBuildInfo(
544 device,
545 CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE,
546 paramValueSize,
547 paramValue,
548 ¶mValueSizeRet);
549
550 EXPECT_EQ(CL_SUCCESS, retVal);
551 EXPECT_EQ(paramValueSizeRet, sizeof(globalVarSize));
552 EXPECT_EQ(globalVarSize, 0u);
553
554 // Set GlobalVariableTotalSize as 1024
555 CreateProgramFromBinary(pContext, pContext->getDevices(), binaryFileName);
556 MockProgram *p = pProgram;
557 ProgramInfo programInfo;
558
559 char constantData[1024] = {};
560 programInfo.globalVariables.initData = constantData;
561 programInfo.globalVariables.size = sizeof(constantData);
562 p->processProgramInfo(programInfo, *pClDevice);
563
564 // get build info once again
565 retVal = pProgram->getBuildInfo(
566 device,
567 CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE,
568 paramValueSize,
569 paramValue,
570 ¶mValueSizeRet);
571 EXPECT_EQ(CL_SUCCESS, retVal);
572 EXPECT_EQ(paramValueSizeRet, sizeof(globalVarSize));
573 if (castToObject<ClDevice>(pClDevice)->areOcl21FeaturesEnabled()) {
574 EXPECT_EQ(globalVarSize, 1024u);
575 } else {
576 EXPECT_EQ(globalVarSize, 0u);
577 }
578 }
579
TEST_F(ProgramFromBinaryTest,givenProgramWhenItIsBeingBuildThenItContainsGraphicsAllocationInKernelInfo)580 TEST_F(ProgramFromBinaryTest, givenProgramWhenItIsBeingBuildThenItContainsGraphicsAllocationInKernelInfo) {
581 pProgram->build(pProgram->getDevices(), nullptr, true);
582 auto kernelInfo = pProgram->getKernelInfo(size_t(0), rootDeviceIndex);
583
584 auto graphicsAllocation = kernelInfo->getGraphicsAllocation();
585 ASSERT_NE(nullptr, graphicsAllocation);
586 EXPECT_TRUE(graphicsAllocation->is32BitAllocation());
587 auto &hwHelper = NEO::HwHelper::get(defaultHwInfo->platform.eRenderCoreFamily);
588 size_t isaPadding = hwHelper.getPaddingForISAAllocation();
589 EXPECT_EQ(graphicsAllocation->getUnderlyingBufferSize(), kernelInfo->heapInfo.KernelHeapSize + isaPadding);
590
591 auto kernelIsa = graphicsAllocation->getUnderlyingBuffer();
592 EXPECT_NE(kernelInfo->heapInfo.pKernelHeap, kernelIsa);
593 EXPECT_EQ(0, memcmp(kernelIsa, kernelInfo->heapInfo.pKernelHeap, kernelInfo->heapInfo.KernelHeapSize));
594 auto rootDeviceIndex = graphicsAllocation->getRootDeviceIndex();
595 EXPECT_EQ(GmmHelper::decanonize(graphicsAllocation->getGpuBaseAddress()), pDevice->getMemoryManager()->getInternalHeapBaseAddress(rootDeviceIndex, graphicsAllocation->isAllocatedInLocalMemoryPool()));
596 }
597
TEST_F(ProgramFromBinaryTest,whenProgramIsBeingRebuildThenOutdatedGlobalBuffersAreFreed)598 TEST_F(ProgramFromBinaryTest, whenProgramIsBeingRebuildThenOutdatedGlobalBuffersAreFreed) {
599 pProgram->build(pProgram->getDevices(), nullptr, true);
600 EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].constantSurface);
601 EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].globalSurface);
602
603 pProgram->buildInfos[pClDevice->getRootDeviceIndex()].constantSurface = new MockGraphicsAllocation();
604 pProgram->processGenBinary(*pClDevice);
605 EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].constantSurface);
606 EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].globalSurface);
607
608 pProgram->buildInfos[pClDevice->getRootDeviceIndex()].globalSurface = new MockGraphicsAllocation();
609 pProgram->processGenBinary(*pClDevice);
610 EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].constantSurface);
611 EXPECT_EQ(nullptr, pProgram->buildInfos[pClDevice->getRootDeviceIndex()].globalSurface);
612 }
613
TEST_F(ProgramFromBinaryTest,givenProgramWhenCleanKernelInfoIsCalledThenKernelAllocationIsFreed)614 TEST_F(ProgramFromBinaryTest, givenProgramWhenCleanKernelInfoIsCalledThenKernelAllocationIsFreed) {
615 pProgram->build(pProgram->getDevices(), nullptr, true);
616 EXPECT_EQ(1u, pProgram->getNumKernels());
617 for (auto i = 0u; i < pProgram->buildInfos.size(); i++) {
618 pProgram->cleanCurrentKernelInfo(i);
619 }
620 EXPECT_EQ(0u, pProgram->getNumKernels());
621 }
622
HWTEST_F(ProgramFromBinaryTest,givenProgramWhenCleanCurrentKernelInfoIsCalledButGpuIsNotYetDoneThenKernelAllocationIsPutOnDeferredFreeListAndCsrRegistersCacheFlush)623 HWTEST_F(ProgramFromBinaryTest, givenProgramWhenCleanCurrentKernelInfoIsCalledButGpuIsNotYetDoneThenKernelAllocationIsPutOnDeferredFreeListAndCsrRegistersCacheFlush) {
624 auto &csr = pDevice->getGpgpuCommandStreamReceiver();
625 EXPECT_TRUE(csr.getTemporaryAllocations().peekIsEmpty());
626 pProgram->build(pProgram->getDevices(), nullptr, true);
627 auto kernelAllocation = pProgram->getKernelInfo(static_cast<size_t>(0u), rootDeviceIndex)->getGraphicsAllocation();
628 kernelAllocation->updateTaskCount(100, csr.getOsContext().getContextId());
629 *csr.getTagAddress() = 0;
630 pProgram->cleanCurrentKernelInfo(rootDeviceIndex);
631 EXPECT_FALSE(csr.getTemporaryAllocations().peekIsEmpty());
632 EXPECT_EQ(csr.getTemporaryAllocations().peekHead(), kernelAllocation);
633 EXPECT_TRUE(this->pDevice->getUltCommandStreamReceiver<FamilyType>().requiresInstructionCacheFlush);
634 }
635
HWTEST_F(ProgramFromBinaryTest,givenIsaAllocationUsedByMultipleCsrsWhenItIsDeletedThenItRegistersCacheFlushInEveryCsrThatUsedIt)636 HWTEST_F(ProgramFromBinaryTest, givenIsaAllocationUsedByMultipleCsrsWhenItIsDeletedThenItRegistersCacheFlushInEveryCsrThatUsedIt) {
637 auto &csr0 = this->pDevice->getUltCommandStreamReceiverFromIndex<FamilyType>(0u);
638 auto &csr1 = this->pDevice->getUltCommandStreamReceiverFromIndex<FamilyType>(1u);
639
640 pProgram->build(pProgram->getDevices(), nullptr, true);
641
642 auto kernelAllocation = pProgram->getKernelInfo(static_cast<size_t>(0u), rootDeviceIndex)->getGraphicsAllocation();
643
644 csr0.makeResident(*kernelAllocation);
645 csr1.makeResident(*kernelAllocation);
646
647 csr0.processResidency(csr0.getResidencyAllocations(), 0u);
648 csr1.processResidency(csr1.getResidencyAllocations(), 0u);
649
650 csr0.makeNonResident(*kernelAllocation);
651 csr1.makeNonResident(*kernelAllocation);
652
653 EXPECT_FALSE(csr0.requiresInstructionCacheFlush);
654 EXPECT_FALSE(csr1.requiresInstructionCacheFlush);
655
656 pProgram->cleanCurrentKernelInfo(rootDeviceIndex);
657 EXPECT_TRUE(csr0.requiresInstructionCacheFlush);
658 EXPECT_TRUE(csr1.requiresInstructionCacheFlush);
659 }
660
TEST_F(ProgramFromSourceTest,GivenSpecificParamatersWhenBuildingProgramThenSuccessOrCorrectErrorCodeIsReturned)661 TEST_F(ProgramFromSourceTest, GivenSpecificParamatersWhenBuildingProgramThenSuccessOrCorrectErrorCodeIsReturned) {
662 KernelBinaryHelper kbHelper(binaryFileName, true);
663 auto device = pPlatform->getClDevice(0);
664
665 CreateProgramWithSource(
666 pContext,
667 sourceFileName);
668
669 // Order of following microtests is important - do not change.
670 // Add new microtests at end.
671
672 auto pMockProgram = pProgram;
673
674 // fail build - another build is already in progress
675 pMockProgram->setBuildStatus(CL_BUILD_IN_PROGRESS);
676 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
677 EXPECT_EQ(CL_INVALID_OPERATION, retVal);
678 pMockProgram->setBuildStatus(CL_BUILD_NONE);
679
680 // fail build - CompilerInterface cannot be obtained
681
682 auto executionEnvironment = device->getExecutionEnvironment();
683 std::unique_ptr<RootDeviceEnvironment> rootDeviceEnvironment = std::make_unique<NoCompilerInterfaceRootDeviceEnvironment>(*executionEnvironment);
684 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
685 auto p2 = std::make_unique<MockProgram>(toClDeviceVector(*device));
686 retVal = p2->build(p2->getDevices(), nullptr, false);
687 EXPECT_EQ(CL_OUT_OF_HOST_MEMORY, retVal);
688 p2.reset(nullptr);
689 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
690
691 // fail build - any build error (here caused by specifying unrecognized option)
692 retVal = pProgram->build(pProgram->getDevices(), "-invalid-option", false);
693 EXPECT_EQ(CL_BUILD_PROGRAM_FAILURE, retVal);
694
695 // fail build - linked code is corrupted and cannot be postprocessed
696 auto p3 = std::make_unique<FailingGenBinaryProgram>(toClDeviceVector(*device));
697 std::string testFile;
698 size_t sourceSize;
699 testFile.append(clFiles);
700 testFile.append("CopyBuffer_simd16.cl"); // source file
701 auto pSourceBuffer = loadDataFromFile(testFile.c_str(), sourceSize);
702 EXPECT_NE(0u, sourceSize);
703 EXPECT_NE(nullptr, pSourceBuffer);
704 p3->sourceCode = pSourceBuffer.get();
705 p3->createdFrom = Program::CreatedFrom::SOURCE;
706 retVal = p3->build(p3->getDevices(), nullptr, false);
707 EXPECT_EQ(CL_INVALID_BINARY, retVal);
708 p3.reset(nullptr);
709
710 // build successfully - build kernel and write it to Kernel Cache
711 pMockProgram->clearOptions();
712 std::string receivedInternalOptions;
713
714 auto debugVars = NEO::getFclDebugVars();
715 debugVars.receivedInternalOptionsOutput = &receivedInternalOptions;
716 gEnvironment->fclPushDebugVars(debugVars);
717 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
718 EXPECT_EQ(CL_SUCCESS, retVal);
719 EXPECT_TRUE(CompilerOptions::contains(receivedInternalOptions, pPlatform->getClDevice(0)->peekCompilerExtensions())) << receivedInternalOptions;
720 gEnvironment->fclPopDebugVars();
721
722 // get build log
723 size_t param_value_size_ret = 0u;
724 retVal = pProgram->getBuildInfo(
725 device,
726 CL_PROGRAM_BUILD_LOG,
727 0,
728 nullptr,
729 ¶m_value_size_ret);
730 EXPECT_EQ(CL_SUCCESS, retVal);
731 EXPECT_NE(param_value_size_ret, 0u);
732
733 // get build log when the log does not exist
734 pMockProgram->clearLog(device->getRootDeviceIndex());
735 retVal = pProgram->getBuildInfo(
736 device,
737 CL_PROGRAM_BUILD_LOG,
738 0,
739 nullptr,
740 ¶m_value_size_ret);
741 EXPECT_EQ(CL_SUCCESS, retVal);
742 EXPECT_NE(param_value_size_ret, 0u);
743
744 // build successfully - build kernel but do not write it to Kernel Cache (kernel is already in the Cache)
745 pMockProgram->setBuildStatus(CL_BUILD_NONE);
746 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
747 EXPECT_EQ(CL_SUCCESS, retVal);
748
749 // build successfully - kernel is already in Kernel Cache, do not build and take it from Cache
750 retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
751 EXPECT_EQ(CL_SUCCESS, retVal);
752
753 // fail build - code to be build does not exist
754 pMockProgram->sourceCode = ""; // set source code as non-existent (invalid)
755 pMockProgram->createdFrom = Program::CreatedFrom::SOURCE;
756 pMockProgram->setBuildStatus(CL_BUILD_NONE);
757 pMockProgram->setCreatedFromBinary(false);
758 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
759 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
760 }
761
TEST_F(ProgramFromSourceTest,GivenDuplicateOptionsWhenCreatingWithSourceThenBuildSucceeds)762 TEST_F(ProgramFromSourceTest, GivenDuplicateOptionsWhenCreatingWithSourceThenBuildSucceeds) {
763 KernelBinaryHelper kbHelper(binaryFileName, false);
764
765 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
766 EXPECT_EQ(CL_SUCCESS, retVal);
767
768 retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::fastRelaxedMath.data(), false);
769 EXPECT_EQ(CL_SUCCESS, retVal);
770
771 retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::fastRelaxedMath.data(), false);
772 EXPECT_EQ(CL_SUCCESS, retVal);
773
774 retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::finiteMathOnly.data(), false);
775 EXPECT_EQ(CL_SUCCESS, retVal);
776
777 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
778 EXPECT_EQ(CL_SUCCESS, retVal);
779 }
780
TEST_F(ProgramFromSourceTest,WhenBuildingProgramThenFeaturesAndExtraExtensionsAreNotAdded)781 TEST_F(ProgramFromSourceTest, WhenBuildingProgramThenFeaturesAndExtraExtensionsAreNotAdded) {
782 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
783 auto pClDevice = pContext->getDevice(0);
784 pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
785
786 auto extensionsOption = static_cast<ClDevice *>(devices[0])->peekCompilerExtensions();
787 auto extensionsWithFeaturesOption = static_cast<ClDevice *>(devices[0])->peekCompilerExtensionsWithFeatures();
788 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsOption)));
789 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsWithFeaturesOption)));
790 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "})));
791
792 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
793 EXPECT_THAT(cip->buildInternalOptions, testing::HasSubstr(extensionsOption));
794 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsWithFeaturesOption)));
795 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "})));
796 }
797
TEST_F(ProgramFromSourceTest,WhenBuildingProgramWithOpenClC20ThenExtraExtensionsAreAdded)798 TEST_F(ProgramFromSourceTest, WhenBuildingProgramWithOpenClC20ThenExtraExtensionsAreAdded) {
799 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
800 auto pClDevice = pContext->getDevice(0);
801 pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
802 auto pProgram = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pClDevice));
803 pProgram->sourceCode = "__kernel mock() {}";
804 pProgram->createdFrom = Program::CreatedFrom::SOURCE;
805
806 MockProgram::initInternalOptionsCalled = 0;
807
808 auto extensionsOption = static_cast<ClDevice *>(devices[0])->peekCompilerExtensions();
809 auto extensionsWithFeaturesOption = static_cast<ClDevice *>(devices[0])->peekCompilerExtensionsWithFeatures();
810 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "})));
811
812 retVal = pProgram->build(pProgram->getDevices(), "-cl-std=CL2.0", false);
813 EXPECT_EQ(CL_SUCCESS, retVal);
814 EXPECT_THAT(cip->buildInternalOptions, testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "}));
815 EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
816 }
817
TEST_F(ProgramFromSourceTest,WhenBuildingProgramWithOpenClC30ThenFeaturesAreAdded)818 TEST_F(ProgramFromSourceTest, WhenBuildingProgramWithOpenClC30ThenFeaturesAreAdded) {
819 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
820 auto pClDevice = pContext->getDevice(0);
821 pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
822 auto pProgram = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pClDevice));
823 pProgram->sourceCode = "__kernel mock() {}";
824 pProgram->createdFrom = Program::CreatedFrom::SOURCE;
825
826 MockProgram::initInternalOptionsCalled = 0;
827
828 auto extensionsOption = static_cast<ClDevice *>(devices[0])->peekCompilerExtensions();
829 auto extensionsWithFeaturesOption = static_cast<ClDevice *>(devices[0])->peekCompilerExtensionsWithFeatures();
830 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsOption)));
831 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsWithFeaturesOption)));
832
833 retVal = pProgram->build(pProgram->getDevices(), "-cl-std=CL3.0", false);
834 EXPECT_EQ(CL_SUCCESS, retVal);
835 EXPECT_THAT(cip->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsOption)));
836 EXPECT_THAT(cip->buildInternalOptions, testing::HasSubstr(extensionsWithFeaturesOption));
837 EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
838 }
839
TEST_F(ProgramFromSourceTest,WhenBuildingProgramWithOpenClC30ThenFeaturesAreAddedOnlyOnce)840 TEST_F(ProgramFromSourceTest, WhenBuildingProgramWithOpenClC30ThenFeaturesAreAddedOnlyOnce) {
841 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
842 auto pClDevice = pContext->getDevice(0);
843 pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
844 auto pProgram = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pClDevice));
845 pProgram->sourceCode = "__kernel mock() {}";
846 pProgram->createdFrom = Program::CreatedFrom::SOURCE;
847
848 retVal = pProgram->build(pProgram->getDevices(), "-cl-std=CL3.0", false);
849 EXPECT_EQ(CL_SUCCESS, retVal);
850 retVal = pProgram->build(pProgram->getDevices(), "-cl-std=CL3.0", false);
851 EXPECT_EQ(CL_SUCCESS, retVal);
852
853 auto extensionsWithFeaturesOption = pClDevice->peekCompilerExtensionsWithFeatures();
854 auto &internalOptions = cip->buildInternalOptions;
855 auto pos = internalOptions.find(extensionsWithFeaturesOption);
856 EXPECT_NE(std::string::npos, pos);
857
858 pos = internalOptions.find(extensionsWithFeaturesOption, pos + 1);
859 EXPECT_EQ(std::string::npos, pos);
860 }
861
TEST_F(ProgramFromSourceTest,WhenCompilingProgramThenFeaturesAndExtraExtensionsAreNotAdded)862 TEST_F(ProgramFromSourceTest, WhenCompilingProgramThenFeaturesAndExtraExtensionsAreNotAdded) {
863 auto pCompilerInterface = new MockCompilerInterfaceCaptureBuildOptions();
864 auto pClDevice = static_cast<ClDevice *>(devices[0]);
865 pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(pCompilerInterface);
866 auto extensionsOption = pClDevice->peekCompilerExtensions();
867 auto extensionsWithFeaturesOption = pClDevice->peekCompilerExtensionsWithFeatures();
868 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsOption)));
869 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsWithFeaturesOption)));
870 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "})));
871
872 MockProgram::initInternalOptionsCalled = 0;
873 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
874 EXPECT_EQ(CL_SUCCESS, retVal);
875 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::HasSubstr(extensionsOption));
876 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsWithFeaturesOption)));
877 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "})));
878 EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
879 }
880
TEST_F(ProgramFromSourceTest,WhenCompilingProgramWithOpenClC20ThenExtraExtensionsAreAdded)881 TEST_F(ProgramFromSourceTest, WhenCompilingProgramWithOpenClC20ThenExtraExtensionsAreAdded) {
882 auto pCompilerInterface = new MockCompilerInterfaceCaptureBuildOptions();
883 auto pClDevice = static_cast<ClDevice *>(devices[0]);
884 pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(pCompilerInterface);
885 auto extensionsOption = pClDevice->peekCompilerExtensions();
886 auto extensionsWithFeaturesOption = pClDevice->peekCompilerExtensionsWithFeatures();
887 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "})));
888
889 MockProgram::initInternalOptionsCalled = 0;
890 retVal = pProgram->compile(pProgram->getDevices(), "-cl-std=CL2.0", 0, nullptr, nullptr);
891 EXPECT_EQ(CL_SUCCESS, retVal);
892 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::HasSubstr(std::string{"+cl_khr_3d_image_writes "}));
893 EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
894 }
895
TEST_F(ProgramFromSourceTest,WhenCompilingProgramWithOpenClC30ThenFeaturesAreAdded)896 TEST_F(ProgramFromSourceTest, WhenCompilingProgramWithOpenClC30ThenFeaturesAreAdded) {
897 auto pCompilerInterface = new MockCompilerInterfaceCaptureBuildOptions();
898 auto pClDevice = pContext->getDevice(0);
899 pClDevice->getExecutionEnvironment()->rootDeviceEnvironments[pClDevice->getRootDeviceIndex()]->compilerInterface.reset(pCompilerInterface);
900 auto pProgram = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pClDevice));
901 pProgram->sourceCode = "__kernel mock() {}";
902 pProgram->createdFrom = Program::CreatedFrom::SOURCE;
903
904 auto extensionsOption = pClDevice->peekCompilerExtensions();
905 auto extensionsWithFeaturesOption = pClDevice->peekCompilerExtensionsWithFeatures();
906 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsOption)));
907 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsWithFeaturesOption)));
908
909 retVal = pProgram->compile(pProgram->getDevices(), "-cl-std=CL3.0", 0, nullptr, nullptr);
910 EXPECT_EQ(CL_SUCCESS, retVal);
911 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::Not(testing::HasSubstr(extensionsOption)));
912 EXPECT_THAT(pCompilerInterface->buildInternalOptions, testing::HasSubstr(extensionsWithFeaturesOption));
913 }
914
915 class Callback {
916 public:
Callback()917 Callback() {
918 this->oldCallback = MemoryManagement::deleteCallback;
919 MemoryManagement::deleteCallback = thisCallback;
920 }
~Callback()921 ~Callback() {
922 MemoryManagement::deleteCallback = this->oldCallback;
923 }
watch(const void * p)924 static void watch(const void *p) {
925 watchList[p] = 0u;
926 }
unwatch(const void * p)927 static void unwatch(const void *p) {
928 EXPECT_GT(watchList[p], 0u);
929 watchList.erase(p);
930 }
931
932 private:
933 void (*oldCallback)(void *);
thisCallback(void * p)934 static void thisCallback(void *p) {
935 if (watchList.find(p) != watchList.end())
936 watchList[p]++;
937 }
938 static std::map<const void *, uint32_t> watchList;
939 };
940
941 std::map<const void *, uint32_t> Callback::watchList;
942
TEST_F(ProgramFromSourceTest,GivenDifferentCommpilerOptionsWhenBuildingProgramThenKernelHashesAreDifferent)943 TEST_F(ProgramFromSourceTest, GivenDifferentCommpilerOptionsWhenBuildingProgramThenKernelHashesAreDifferent) {
944 KernelBinaryHelper kbHelper(binaryFileName, true);
945
946 auto rootDeviceIndex = pContext->getDevice(0)->getRootDeviceIndex();
947
948 CreateProgramWithSource(
949 pContext,
950 sourceFileName);
951
952 Callback callback;
953
954 retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
955 EXPECT_EQ(CL_SUCCESS, retVal);
956 auto hash1 = pProgram->getCachedFileName();
957 auto kernel1 = pProgram->getKernelInfo("CopyBuffer", rootDeviceIndex);
958 Callback::watch(kernel1);
959 EXPECT_NE(nullptr, kernel1);
960
961 retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::fastRelaxedMath.data(), true);
962 EXPECT_EQ(CL_SUCCESS, retVal);
963 auto hash2 = pProgram->getCachedFileName();
964 auto kernel2 = pProgram->getKernelInfo("CopyBuffer", rootDeviceIndex);
965 EXPECT_NE(nullptr, kernel2);
966 EXPECT_NE(hash1, hash2);
967 Callback::unwatch(kernel1);
968 Callback::watch(kernel2);
969
970 retVal = pProgram->build(pProgram->getDevices(), CompilerOptions::finiteMathOnly.data(), true);
971 EXPECT_EQ(CL_SUCCESS, retVal);
972 auto hash3 = pProgram->getCachedFileName();
973 auto kernel3 = pProgram->getKernelInfo("CopyBuffer", rootDeviceIndex);
974 EXPECT_NE(nullptr, kernel3);
975 EXPECT_NE(hash1, hash3);
976 EXPECT_NE(hash2, hash3);
977 Callback::unwatch(kernel2);
978 Callback::watch(kernel3);
979
980 pProgram->createdFrom = NEO::Program::CreatedFrom::BINARY;
981 pProgram->setIrBinary(new char[16], true);
982 pProgram->setIrBinarySize(16, true);
983 retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
984 EXPECT_EQ(CL_SUCCESS, retVal);
985 auto hash4 = pProgram->getCachedFileName();
986 auto kernel4 = pProgram->getKernelInfo("CopyBuffer", rootDeviceIndex);
987 EXPECT_NE(nullptr, kernel4);
988 EXPECT_EQ(hash3, hash4);
989 Callback::unwatch(kernel3);
990 Callback::watch(kernel4);
991
992 pProgram->createdFrom = NEO::Program::CreatedFrom::SOURCE;
993 retVal = pProgram->build(pProgram->getDevices(), nullptr, true);
994 EXPECT_EQ(CL_SUCCESS, retVal);
995 auto hash5 = pProgram->getCachedFileName();
996 auto kernel5 = pProgram->getKernelInfo("CopyBuffer", rootDeviceIndex);
997 EXPECT_NE(nullptr, kernel5);
998 EXPECT_EQ(hash1, hash5);
999 Callback::unwatch(kernel4);
1000 }
1001
TEST_F(ProgramFromSourceTest,GivenEmptyProgramWhenCreatingProgramThenInvalidValueErrorIsReturned)1002 TEST_F(ProgramFromSourceTest, GivenEmptyProgramWhenCreatingProgramThenInvalidValueErrorIsReturned) {
1003 auto p = Program::create(pContext, 0, nullptr, nullptr, retVal);
1004 EXPECT_EQ(CL_INVALID_VALUE, retVal);
1005 EXPECT_EQ(nullptr, p);
1006 delete p;
1007 }
1008
TEST_F(ProgramFromSourceTest,GivenSpecificParamatersWhenCompilingProgramThenSuccessOrCorrectErrorCodeIsReturned)1009 TEST_F(ProgramFromSourceTest, GivenSpecificParamatersWhenCompilingProgramThenSuccessOrCorrectErrorCodeIsReturned) {
1010 CreateProgramWithSource(
1011 pContext,
1012 sourceFileName);
1013
1014 cl_program inputHeaders;
1015 const char *headerIncludeNames = "";
1016 cl_program nullprogram = nullptr;
1017 cl_program invprogram = (cl_program)pContext;
1018
1019 // Order of following microtests is important - do not change.
1020 // Add new microtests at end.
1021
1022 // invalid compile parameters: combinations of numInputHeaders==0 & inputHeaders & headerIncludeNames
1023 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, &inputHeaders, nullptr);
1024 EXPECT_EQ(CL_INVALID_VALUE, retVal);
1025
1026 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, &headerIncludeNames);
1027 EXPECT_EQ(CL_INVALID_VALUE, retVal);
1028
1029 // invalid compile parameters: combinations of numInputHeaders!=0 & inputHeaders & headerIncludeNames
1030 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 1, &inputHeaders, nullptr);
1031 EXPECT_EQ(CL_INVALID_VALUE, retVal);
1032
1033 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 1, nullptr, &headerIncludeNames);
1034 EXPECT_EQ(CL_INVALID_VALUE, retVal);
1035
1036 // fail compilation - another compilation is already in progress
1037 pProgram->setBuildStatus(CL_BUILD_IN_PROGRESS);
1038 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
1039 EXPECT_EQ(CL_INVALID_OPERATION, retVal);
1040 pProgram->setBuildStatus(CL_BUILD_NONE);
1041
1042 // invalid compile parameters: invalid header Program object==nullptr
1043 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 1, &nullprogram, &headerIncludeNames);
1044 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
1045
1046 // invalid compile parameters: invalid header Program object==non Program object
1047 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 1, &invprogram, &headerIncludeNames);
1048 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
1049
1050 // compile successfully kernel with header
1051 std::string testFile;
1052 size_t sourceSize;
1053 MockProgram *p3; // header Program object
1054 testFile.append(clFiles);
1055 testFile.append("CopyBuffer_simd16.cl"); // header source file
1056 auto pSourceBuffer = loadDataFromFile(testFile.c_str(), sourceSize);
1057 EXPECT_NE(0u, sourceSize);
1058 EXPECT_NE(nullptr, pSourceBuffer);
1059 const char *sources[1] = {pSourceBuffer.get()};
1060 p3 = Program::create<MockProgram>(pContext, 1, sources, &sourceSize, retVal);
1061 EXPECT_EQ(CL_SUCCESS, retVal);
1062 EXPECT_NE(nullptr, p3);
1063 inputHeaders = p3;
1064 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 1, &inputHeaders, &headerIncludeNames);
1065 EXPECT_EQ(CL_SUCCESS, retVal);
1066
1067 // fail compilation of kernel with header - header is invalid
1068 p3->sourceCode = ""; // set header source code as non-existent (invalid)
1069 retVal = p3->compile(p3->getDevices(), nullptr, 1, &inputHeaders, &headerIncludeNames);
1070 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
1071 delete p3;
1072
1073 // fail compilation - CompilerInterface cannot be obtained
1074 auto device = pContext->getDevice(0);
1075 auto executionEnvironment = device->getExecutionEnvironment();
1076 std::unique_ptr<RootDeviceEnvironment> rootDeviceEnvironment = std::make_unique<NoCompilerInterfaceRootDeviceEnvironment>(*executionEnvironment);
1077 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
1078 auto p2 = std::make_unique<MockProgram>(toClDeviceVector(*device));
1079 retVal = p2->compile(p2->getDevices(), nullptr, 0, nullptr, nullptr);
1080 EXPECT_EQ(CL_OUT_OF_HOST_MEMORY, retVal);
1081 p2.reset(nullptr);
1082 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
1083
1084 // fail compilation - any compilation error (here caused by specifying unrecognized option)
1085 retVal = pProgram->compile(pProgram->getDevices(), "-invalid-option", 0, nullptr, nullptr);
1086 EXPECT_EQ(CL_COMPILE_PROGRAM_FAILURE, retVal);
1087
1088 // compile successfully
1089 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
1090 EXPECT_EQ(CL_SUCCESS, retVal);
1091 }
1092
TEST_F(ProgramFromSourceTest,GivenFlagsWhenCompilingProgramThenBuildOptionsHaveBeenApplied)1093 TEST_F(ProgramFromSourceTest, GivenFlagsWhenCompilingProgramThenBuildOptionsHaveBeenApplied) {
1094 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
1095 auto pDevice = pContext->getDevice(0);
1096 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
1097 auto program = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pDevice));
1098 program->sourceCode = "__kernel mock() {}";
1099
1100 // Ask to build created program without NEO::CompilerOptions::gtpinRera and NEO::CompilerOptions::greaterThan4gbBuffersRequired flags.
1101 cl_int retVal = program->compile(pProgram->getDevices(), CompilerOptions::fastRelaxedMath.data(), 0, nullptr, nullptr);
1102 EXPECT_EQ(CL_SUCCESS, retVal);
1103
1104 // Check build options that were applied
1105 EXPECT_TRUE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::fastRelaxedMath)) << cip->buildOptions;
1106 EXPECT_FALSE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::gtpinRera)) << cip->buildInternalOptions;
1107 if (!pDevice->areSharedSystemAllocationsAllowed()) {
1108 EXPECT_FALSE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::greaterThan4gbBuffersRequired)) << cip->buildInternalOptions;
1109 }
1110 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, pPlatform->getClDevice(0)->peekCompilerExtensions())) << cip->buildInternalOptions;
1111
1112 // Ask to build created program with NEO::CompilerOptions::gtpinRera and NEO::CompilerOptions::greaterThan4gbBuffersRequired flags.
1113 cip->buildOptions.clear();
1114 cip->buildInternalOptions.clear();
1115 auto options = CompilerOptions::concatenate(CompilerOptions::greaterThan4gbBuffersRequired, CompilerOptions::gtpinRera, CompilerOptions::finiteMathOnly);
1116 retVal = program->compile(pProgram->getDevices(), options.c_str(),
1117 0, nullptr, nullptr);
1118 EXPECT_EQ(CL_SUCCESS, retVal);
1119
1120 // Check build options that were applied
1121 EXPECT_FALSE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::fastRelaxedMath)) << cip->buildOptions;
1122 EXPECT_TRUE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::finiteMathOnly)) << cip->buildOptions;
1123 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::gtpinRera)) << cip->buildInternalOptions;
1124 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::greaterThan4gbBuffersRequired)) << cip->buildInternalOptions;
1125 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, pPlatform->getClDevice(0)->peekCompilerExtensions())) << cip->buildInternalOptions;
1126 }
1127
TEST_F(ProgramTests,GivenFlagsWhenLinkingProgramThenBuildOptionsHaveBeenApplied)1128 TEST_F(ProgramTests, GivenFlagsWhenLinkingProgramThenBuildOptionsHaveBeenApplied) {
1129 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
1130 auto pProgram = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pClDevice));
1131 pProgram->sourceCode = "__kernel mock() {}";
1132 pProgram->createdFrom = Program::CreatedFrom::SOURCE;
1133 MockProgram::initInternalOptionsCalled = 0;
1134
1135 cl_program program = pProgram.get();
1136
1137 // compile successfully a kernel to be linked later
1138 cl_int retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
1139 EXPECT_EQ(CL_SUCCESS, retVal);
1140 EXPECT_EQ(1, MockProgram::initInternalOptionsCalled);
1141
1142 // Ask to link created program with NEO::CompilerOptions::gtpinRera and NEO::CompilerOptions::greaterThan4gbBuffersRequired flags.
1143 auto options = CompilerOptions::concatenate(CompilerOptions::greaterThan4gbBuffersRequired, CompilerOptions::gtpinRera, CompilerOptions::finiteMathOnly);
1144
1145 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
1146
1147 retVal = pProgram->link(pProgram->getDevices(), options.c_str(), 1, &program);
1148 EXPECT_EQ(CL_SUCCESS, retVal);
1149 EXPECT_EQ(2, MockProgram::initInternalOptionsCalled);
1150
1151 // Check build options that were applied
1152 EXPECT_FALSE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::fastRelaxedMath)) << cip->buildOptions;
1153 EXPECT_TRUE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::finiteMathOnly)) << cip->buildOptions;
1154 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::gtpinRera)) << cip->buildInternalOptions;
1155 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::greaterThan4gbBuffersRequired)) << cip->buildInternalOptions;
1156 }
1157
TEST_F(ProgramFromSourceTest,GivenAdvancedOptionsWhenCreatingProgramThenSuccessIsReturned)1158 TEST_F(ProgramFromSourceTest, GivenAdvancedOptionsWhenCreatingProgramThenSuccessIsReturned) {
1159 std::string testFile;
1160 size_t sourceSize = 0;
1161
1162 Program *p;
1163 testFile.append(clFiles);
1164 testFile.append("CopyBuffer_simd16.cl");
1165 auto pSourceBuffer = loadDataFromFile(testFile.c_str(), sourceSize);
1166 const char *sources[1] = {pSourceBuffer.get()};
1167 EXPECT_NE(nullptr, pSourceBuffer);
1168
1169 //According to spec: If lengths is NULL, all strings in the strings argument are considered null-terminated.
1170 p = Program::create(pContext, 1, sources, nullptr, retVal);
1171 EXPECT_EQ(CL_SUCCESS, retVal);
1172 EXPECT_NE(nullptr, p);
1173 delete p;
1174
1175 //According to spec: If an element in lengths is zero, its accompanying string is null-terminated.
1176 p = Program::create(pContext, 1, sources, &sourceSize, retVal);
1177 EXPECT_EQ(CL_SUCCESS, retVal);
1178 EXPECT_NE(nullptr, p);
1179 delete p;
1180
1181 std::stringstream dataStream(pSourceBuffer.get());
1182 std::string line;
1183 std::vector<const char *> lines;
1184 while (std::getline(dataStream, line, '\n')) {
1185 char *ptr = new char[line.length() + 1]();
1186 strcpy_s(ptr, line.length() + 1, line.c_str());
1187 lines.push_back(ptr);
1188 }
1189
1190 // Work on array of strings
1191 p = Program::create(pContext, 1, &lines[0], nullptr, retVal);
1192 EXPECT_EQ(CL_SUCCESS, retVal);
1193 EXPECT_NE(nullptr, p);
1194 delete p;
1195
1196 std::vector<size_t> sizes;
1197 for (auto ptr : lines)
1198 sizes.push_back(strlen(ptr));
1199 sizes[sizes.size() / 2] = 0;
1200
1201 p = Program::create(pContext, (cl_uint)sizes.size(), &lines[0], &sizes[0], retVal);
1202 EXPECT_EQ(CL_SUCCESS, retVal);
1203 EXPECT_NE(nullptr, p);
1204 delete p;
1205
1206 for (auto ptr : lines)
1207 delete[] ptr;
1208 }
1209
TEST_F(ProgramFromSourceTest,GivenSpecificParamatersWhenLinkingProgramThenSuccessOrCorrectErrorCodeIsReturned)1210 TEST_F(ProgramFromSourceTest, GivenSpecificParamatersWhenLinkingProgramThenSuccessOrCorrectErrorCodeIsReturned) {
1211 CreateProgramWithSource(
1212 pContext,
1213 sourceFileName);
1214
1215 cl_program program = pProgram;
1216 cl_program nullprogram = nullptr;
1217 cl_program invprogram = (cl_program)pContext;
1218
1219 // Order of following microtests is important - do not change.
1220 // Add new microtests at end.
1221
1222 // invalid link parameters: combinations of numInputPrograms & inputPrograms
1223 retVal = pProgram->link(pProgram->getDevices(), nullptr, 0, &program);
1224 EXPECT_EQ(CL_INVALID_VALUE, retVal);
1225
1226 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, nullptr);
1227 EXPECT_EQ(CL_INVALID_VALUE, retVal);
1228
1229 // fail linking - another linking is already in progress
1230 pProgram->setBuildStatus(CL_BUILD_IN_PROGRESS);
1231 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program);
1232 EXPECT_EQ(CL_INVALID_OPERATION, retVal);
1233 pProgram->setBuildStatus(CL_BUILD_NONE);
1234
1235 // invalid link parameters: invalid Program object==nullptr
1236 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &nullprogram);
1237 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
1238
1239 // invalid link parameters: invalid Program object==non Program object
1240 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &invprogram);
1241 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
1242
1243 // compile successfully a kernel to be linked later
1244 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
1245 EXPECT_EQ(CL_SUCCESS, retVal);
1246
1247 // fail linking - code to be linked does not exist
1248 bool isSpirvTmp = pProgram->getIsSpirV();
1249 char *pIrBin = pProgram->irBinary.get();
1250 pProgram->irBinary.release();
1251 size_t irBinSize = pProgram->irBinarySize;
1252 pProgram->setIrBinary(nullptr, false);
1253 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program);
1254 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
1255 pProgram->setIrBinary(pIrBin, isSpirvTmp);
1256
1257 // fail linking - size of code to be linked is == 0
1258 pProgram->setIrBinarySize(0, isSpirvTmp);
1259 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program);
1260 EXPECT_EQ(CL_INVALID_PROGRAM, retVal);
1261 pProgram->setIrBinarySize(irBinSize, isSpirvTmp);
1262
1263 // fail linking - any link error (here caused by specifying unrecognized option)
1264 retVal = pProgram->link(pProgram->getDevices(), "-invalid-option", 1, &program);
1265 EXPECT_EQ(CL_LINK_PROGRAM_FAILURE, retVal);
1266
1267 // fail linking - linked code is corrupted and cannot be postprocessed
1268 auto p2 = std::make_unique<FailingGenBinaryProgram>(pProgram->getDevices());
1269 retVal = p2->link(p2->getDevices(), nullptr, 1, &program);
1270 EXPECT_EQ(CL_INVALID_BINARY, retVal);
1271 p2.reset(nullptr);
1272
1273 // link successfully
1274 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &program);
1275 EXPECT_EQ(CL_SUCCESS, retVal);
1276 }
1277
TEST_F(ProgramFromSourceTest,GivenInvalidOptionsWhenCreatingLibraryThenCorrectErrorIsReturned)1278 TEST_F(ProgramFromSourceTest, GivenInvalidOptionsWhenCreatingLibraryThenCorrectErrorIsReturned) {
1279 cl_program program = pProgram;
1280
1281 // Order of following microtests is important - do not change.
1282 // Add new microtests at end.
1283
1284 // compile successfully a kernel to be later used to create library
1285 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
1286 EXPECT_EQ(CL_SUCCESS, retVal);
1287
1288 // create library successfully
1289 retVal = pProgram->link(pProgram->getDevices(), CompilerOptions::createLibrary.data(), 1, &program);
1290 EXPECT_EQ(CL_SUCCESS, retVal);
1291
1292 // fail library creation - any link error (here caused by specifying unrecognized option)
1293 retVal = pProgram->link(pProgram->getDevices(), CompilerOptions::concatenate(CompilerOptions::createLibrary, "-invalid-option").c_str(), 1, &program);
1294 EXPECT_EQ(CL_LINK_PROGRAM_FAILURE, retVal);
1295
1296 auto device = pContext->getDevice(0);
1297 auto executionEnvironment = device->getExecutionEnvironment();
1298 std::unique_ptr<RootDeviceEnvironment> rootDeviceEnvironment = std::make_unique<NoCompilerInterfaceRootDeviceEnvironment>(*executionEnvironment);
1299 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
1300 auto failingProgram = std::make_unique<MockProgram>(toClDeviceVector(*device));
1301
1302 // fail library creation - CompilerInterface cannot be obtained
1303 retVal = failingProgram->link(failingProgram->getDevices(), CompilerOptions::createLibrary.data(), 1, &program);
1304 EXPECT_EQ(CL_OUT_OF_HOST_MEMORY, retVal);
1305 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[device->getRootDeviceIndex()]);
1306 }
1307
1308 class PatchTokenFromBinaryTest : public ProgramSimpleFixture {
1309 public:
SetUp()1310 void SetUp() override {
1311 ProgramSimpleFixture::SetUp();
1312 }
TearDown()1313 void TearDown() override {
1314 ProgramSimpleFixture::TearDown();
1315 }
1316 };
1317 using PatchTokenTests = Test<PatchTokenFromBinaryTest>;
1318
1319 template <typename FamilyType>
1320 class CommandStreamReceiverMock : public UltCommandStreamReceiver<FamilyType> {
1321 using BaseClass = UltCommandStreamReceiver<FamilyType>;
1322 using BaseClass::BaseClass;
1323
1324 public:
makeResident(GraphicsAllocation & graphicsAllocation)1325 void makeResident(GraphicsAllocation &graphicsAllocation) override {
1326 residency[graphicsAllocation.getUnderlyingBuffer()] = graphicsAllocation.getUnderlyingBufferSize();
1327 CommandStreamReceiver::makeResident(graphicsAllocation);
1328 }
1329
makeNonResident(GraphicsAllocation & graphicsAllocation)1330 void makeNonResident(GraphicsAllocation &graphicsAllocation) override {
1331 residency.erase(graphicsAllocation.getUnderlyingBuffer());
1332 CommandStreamReceiver::makeNonResident(graphicsAllocation);
1333 }
1334
1335 std::map<const void *, size_t> residency;
1336 };
1337
HWTEST_F(PatchTokenTests,givenKernelRequiringConstantAllocationWhenMakeResidentIsCalledThenConstantAllocationIsMadeResident)1338 HWTEST_F(PatchTokenTests, givenKernelRequiringConstantAllocationWhenMakeResidentIsCalledThenConstantAllocationIsMadeResident) {
1339 CreateProgramFromBinary(pContext, pContext->getDevices(), "test_constant_memory");
1340
1341 ASSERT_NE(nullptr, pProgram);
1342 retVal = pProgram->build(
1343 pProgram->getDevices(),
1344 nullptr,
1345 false);
1346
1347 ASSERT_EQ(CL_SUCCESS, retVal);
1348
1349 auto pKernelInfo = pProgram->getKernelInfo("test", rootDeviceIndex);
1350
1351 ASSERT_NE(nullptr, pProgram->getConstantSurface(pClDevice->getRootDeviceIndex()));
1352
1353 uint32_t expected_values[] = {0xabcd5432u, 0xaabb5533u};
1354 uint32_t *constBuff = reinterpret_cast<uint32_t *>(pProgram->getConstantSurface(pClDevice->getRootDeviceIndex())->getUnderlyingBuffer());
1355 EXPECT_EQ(expected_values[0], constBuff[0]);
1356 EXPECT_EQ(expected_values[1], constBuff[1]);
1357
1358 std::unique_ptr<Kernel> pKernel(Kernel::create(pProgram, *pKernelInfo, *pClDevice, &retVal));
1359
1360 ASSERT_EQ(CL_SUCCESS, retVal);
1361 ASSERT_NE(nullptr, pKernel);
1362
1363 auto pCommandStreamReceiver = new CommandStreamReceiverMock<FamilyType>(*pDevice->executionEnvironment, pDevice->getRootDeviceIndex(), pDevice->getDeviceBitfield());
1364 ASSERT_NE(nullptr, pCommandStreamReceiver);
1365
1366 pDevice->resetCommandStreamReceiver(pCommandStreamReceiver);
1367 pCommandStreamReceiver->residency.clear();
1368
1369 pKernel->makeResident(*pCommandStreamReceiver);
1370 EXPECT_EQ(2u, pCommandStreamReceiver->residency.size());
1371
1372 auto &residencyVector = pCommandStreamReceiver->getResidencyAllocations();
1373
1374 //we expect kernel ISA here and constant allocation
1375 auto kernelIsa = pKernel->getKernelInfo().getGraphicsAllocation();
1376 auto constantAllocation = pProgram->getConstantSurface(pDevice->getRootDeviceIndex());
1377
1378 auto element = std::find(residencyVector.begin(), residencyVector.end(), kernelIsa);
1379 EXPECT_NE(residencyVector.end(), element);
1380 element = std::find(residencyVector.begin(), residencyVector.end(), constantAllocation);
1381 EXPECT_NE(residencyVector.end(), element);
1382
1383 auto crossThreadData = pKernel->getCrossThreadData();
1384 uint32_t *constBuffGpuAddr = reinterpret_cast<uint32_t *>(pProgram->getConstantSurface(pContext->getDevice(0)->getRootDeviceIndex())->getGpuAddressToPatch());
1385 uintptr_t *pDst = reinterpret_cast<uintptr_t *>(crossThreadData + pKernelInfo->kernelDescriptor.payloadMappings.implicitArgs.globalConstantsSurfaceAddress.stateless);
1386
1387 EXPECT_EQ(*pDst, reinterpret_cast<uintptr_t>(constBuffGpuAddr));
1388
1389 pCommandStreamReceiver->makeSurfacePackNonResident(pCommandStreamReceiver->getResidencyAllocations());
1390 EXPECT_EQ(0u, pCommandStreamReceiver->residency.size());
1391
1392 std::vector<Surface *> surfaces;
1393 pKernel->getResidency(surfaces);
1394 EXPECT_EQ(2u, surfaces.size());
1395
1396 for (Surface *surface : surfaces) {
1397 delete surface;
1398 }
1399 }
1400
TEST_F(PatchTokenTests,WhenBuildingProgramThenGwsIsSet)1401 TEST_F(PatchTokenTests, WhenBuildingProgramThenGwsIsSet) {
1402 CreateProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
1403
1404 ASSERT_NE(nullptr, pProgram);
1405 retVal = pProgram->build(
1406 pProgram->getDevices(),
1407 nullptr,
1408 false);
1409
1410 ASSERT_EQ(CL_SUCCESS, retVal);
1411
1412 auto pKernelInfo = pProgram->getKernelInfo("test", rootDeviceIndex);
1413
1414 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.globalWorkSize[0]);
1415 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.globalWorkSize[1]);
1416 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.globalWorkSize[2]);
1417 }
1418
TEST_F(PatchTokenTests,WhenBuildingProgramThenLwsIsSet)1419 TEST_F(PatchTokenTests, WhenBuildingProgramThenLwsIsSet) {
1420 CreateProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
1421
1422 ASSERT_NE(nullptr, pProgram);
1423 retVal = pProgram->build(
1424 pProgram->getDevices(),
1425 nullptr,
1426 false);
1427
1428 ASSERT_EQ(CL_SUCCESS, retVal);
1429
1430 auto pKernelInfo = pProgram->getKernelInfo("test", rootDeviceIndex);
1431
1432 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize[0]);
1433 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize[1]);
1434 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize[2]);
1435
1436 pKernelInfo = pProgram->getKernelInfo("test_get_local_size", rootDeviceIndex);
1437
1438 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize[0]);
1439 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize[1]);
1440 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize[2]);
1441 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize2[0]);
1442 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize2[1]);
1443 ASSERT_NE(static_cast<uint32_t>(-1), pKernelInfo->kernelDescriptor.payloadMappings.dispatchTraits.localWorkSize2[2]);
1444 }
1445
TEST_F(PatchTokenTests,WhenBuildingProgramThenConstantKernelArgsAreAvailable)1446 TEST_F(PatchTokenTests, WhenBuildingProgramThenConstantKernelArgsAreAvailable) {
1447 // PATCH_TOKEN_STATELESS_CONSTANT_MEMORY_OBJECT_KERNEL_ARGUMENT
1448
1449 CreateProgramFromBinary(pContext, pContext->getDevices(), "test_basic_constant");
1450
1451 ASSERT_NE(nullptr, pProgram);
1452 retVal = pProgram->build(
1453 pProgram->getDevices(),
1454 nullptr,
1455 false);
1456
1457 EXPECT_EQ(CL_SUCCESS, retVal);
1458
1459 auto pKernelInfo = pProgram->getKernelInfo("constant_kernel", rootDeviceIndex);
1460 ASSERT_NE(nullptr, pKernelInfo);
1461
1462 auto pKernel = Kernel::create(
1463 pProgram,
1464 *pKernelInfo,
1465 *pClDevice,
1466 &retVal);
1467
1468 ASSERT_EQ(CL_SUCCESS, retVal);
1469 ASSERT_NE(nullptr, pKernel);
1470
1471 uint32_t numArgs;
1472 retVal = pKernel->getInfo(CL_KERNEL_NUM_ARGS, sizeof(numArgs), &numArgs, nullptr);
1473 EXPECT_EQ(CL_SUCCESS, retVal);
1474 EXPECT_EQ(3u, numArgs);
1475
1476 uint32_t sizeOfPtr = sizeof(void *);
1477 EXPECT_EQ(pKernelInfo->getArgDescriptorAt(0).as<ArgDescPointer>().pointerSize, sizeOfPtr);
1478 EXPECT_EQ(pKernelInfo->getArgDescriptorAt(1).as<ArgDescPointer>().pointerSize, sizeOfPtr);
1479
1480 delete pKernel;
1481 }
1482
TEST_F(PatchTokenTests,GivenVmeKernelWhenBuildingKernelThenArgAvailable)1483 TEST_F(PatchTokenTests, GivenVmeKernelWhenBuildingKernelThenArgAvailable) {
1484 if (!pDevice->getHardwareInfo().capabilityTable.supportsVme) {
1485 GTEST_SKIP();
1486 }
1487 // PATCH_TOKEN_INLINE_VME_SAMPLER_INFO token indicates a VME kernel.
1488
1489 CreateProgramFromBinary(pContext, pContext->getDevices(), "vme_kernels");
1490
1491 ASSERT_NE(nullptr, pProgram);
1492 retVal = pProgram->build(
1493 pProgram->getDevices(),
1494 nullptr,
1495 false);
1496
1497 EXPECT_EQ(CL_SUCCESS, retVal);
1498
1499 auto pKernelInfo = pProgram->getKernelInfo("device_side_block_motion_estimate_intel", rootDeviceIndex);
1500 ASSERT_NE(nullptr, pKernelInfo);
1501 EXPECT_EQ(true, pKernelInfo->kernelDescriptor.kernelAttributes.flags.usesVme);
1502
1503 auto pKernel = Kernel::create(
1504 pProgram,
1505 *pKernelInfo,
1506 *pClDevice,
1507 &retVal);
1508
1509 ASSERT_NE(nullptr, pKernel);
1510
1511 delete pKernel;
1512 }
1513
1514 class ProgramPatchTokenFromBinaryTest : public ProgramSimpleFixture {
1515 public:
SetUp()1516 void SetUp() override {
1517 ProgramSimpleFixture::SetUp();
1518 }
TearDown()1519 void TearDown() override {
1520 ProgramSimpleFixture::TearDown();
1521 }
1522 };
1523 typedef Test<ProgramPatchTokenFromBinaryTest> ProgramPatchTokenTests;
1524
TEST(ProgramFromBinaryTests,givenBinaryWithInvalidICBEThenErrorIsReturned)1525 TEST(ProgramFromBinaryTests, givenBinaryWithInvalidICBEThenErrorIsReturned) {
1526 cl_int retVal = CL_INVALID_BINARY;
1527
1528 SProgramBinaryHeader binHeader;
1529 memset(&binHeader, 0, sizeof(binHeader));
1530 binHeader.Magic = iOpenCL::MAGIC_CL;
1531 binHeader.Version = iOpenCL::CURRENT_ICBE_VERSION - 3;
1532 binHeader.Device = defaultHwInfo->platform.eRenderCoreFamily;
1533 binHeader.GPUPointerSizeInBytes = 8;
1534 binHeader.NumberOfKernels = 0;
1535 binHeader.SteppingId = 0;
1536 binHeader.PatchListSize = 0;
1537 size_t binSize = sizeof(SProgramBinaryHeader);
1538
1539 {
1540 const unsigned char *binaries[1] = {reinterpret_cast<const unsigned char *>(&binHeader)};
1541 MockContext context;
1542
1543 std::unique_ptr<Program> pProgram(Program::create<Program>(&context, context.getDevices(), &binSize, binaries, nullptr, retVal));
1544 EXPECT_EQ(nullptr, pProgram.get());
1545 EXPECT_EQ(CL_INVALID_BINARY, retVal);
1546 }
1547
1548 {
1549 // whatever method we choose CL_INVALID_BINARY is always returned
1550 auto device = std::make_unique<MockClDevice>(MockDevice::createWithNewExecutionEnvironment<MockDevice>(nullptr, mockRootDeviceIndex));
1551 std::unique_ptr<Program> pProgram(Program::createBuiltInFromGenBinary(nullptr, toClDeviceVector(*device), &binHeader, binSize, &retVal));
1552 ASSERT_NE(nullptr, pProgram.get());
1553 EXPECT_EQ(CL_SUCCESS, retVal);
1554
1555 retVal = pProgram->processGenBinary(*device);
1556 EXPECT_EQ(CL_INVALID_BINARY, retVal);
1557 }
1558 }
1559
TEST(ProgramFromBinaryTests,givenEmptyProgramThenErrorIsReturned)1560 TEST(ProgramFromBinaryTests, givenEmptyProgramThenErrorIsReturned) {
1561 cl_int retVal = CL_INVALID_BINARY;
1562
1563 SProgramBinaryHeader binHeader;
1564 memset(&binHeader, 0, sizeof(binHeader));
1565 binHeader.Magic = iOpenCL::MAGIC_CL;
1566 binHeader.Version = iOpenCL::CURRENT_ICBE_VERSION;
1567 binHeader.Device = defaultHwInfo->platform.eRenderCoreFamily;
1568 binHeader.GPUPointerSizeInBytes = 8;
1569 binHeader.NumberOfKernels = 0;
1570 binHeader.SteppingId = 0;
1571 binHeader.PatchListSize = 0;
1572 size_t binSize = sizeof(SProgramBinaryHeader);
1573
1574 auto device = std::make_unique<MockClDevice>(MockDevice::createWithNewExecutionEnvironment<MockDevice>(nullptr, mockRootDeviceIndex));
1575 std::unique_ptr<MockProgram> pProgram(MockProgram::createBuiltInFromGenBinary<MockProgram>(nullptr, toClDeviceVector(*device), &binHeader, binSize, &retVal));
1576 ASSERT_NE(nullptr, pProgram.get());
1577 EXPECT_EQ(CL_SUCCESS, retVal);
1578
1579 auto rootDeviceIndex = mockRootDeviceIndex;
1580 pProgram->buildInfos[rootDeviceIndex].unpackedDeviceBinary.reset(nullptr);
1581 retVal = pProgram->processGenBinary(*device);
1582 EXPECT_EQ(CL_INVALID_BINARY, retVal);
1583 }
1584
1585 using ProgramWithDebugSymbolsTests = Test<ProgramSimpleFixture>;
1586
TEST_F(ProgramWithDebugSymbolsTests,GivenProgramCreatedWithDashGOptionWhenGettingProgramBinariesThenDebugDataIsIncluded)1587 TEST_F(ProgramWithDebugSymbolsTests, GivenProgramCreatedWithDashGOptionWhenGettingProgramBinariesThenDebugDataIsIncluded) {
1588 CreateProgramFromBinary(pContext, pContext->getDevices(), "CopyBuffer_simd16", "-g");
1589
1590 ASSERT_NE(nullptr, pProgram);
1591
1592 retVal = pProgram->build(
1593 pProgram->getDevices(),
1594 "-g",
1595 false);
1596 EXPECT_EQ(CL_SUCCESS, retVal);
1597
1598 size_t paramValueSize = sizeof(size_t);
1599 size_t paramValueSizeRet = 0;
1600 size_t size = 0;
1601
1602 pProgram->buildInfos[rootDeviceIndex].packedDeviceBinary.reset();
1603 pProgram->buildInfos[rootDeviceIndex].packedDeviceBinarySize = 0U;
1604
1605 retVal = pProgram->packDeviceBinary(*pClDevice);
1606
1607 retVal = pProgram->getInfo(
1608 CL_PROGRAM_BINARY_SIZES,
1609 paramValueSize,
1610 &size,
1611 nullptr);
1612
1613 EXPECT_EQ(CL_SUCCESS, retVal);
1614
1615 auto testBinary = std::make_unique<char[]>(size);
1616
1617 retVal = pProgram->getInfo(
1618 CL_PROGRAM_BINARIES,
1619 paramValueSize,
1620 &testBinary,
1621 ¶mValueSizeRet);
1622
1623 EXPECT_EQ(CL_SUCCESS, retVal);
1624
1625 ArrayRef<const uint8_t> archive(reinterpret_cast<const uint8_t *>(testBinary.get()), size);
1626 auto productAbbreviation = hardwarePrefix[pDevice->getHardwareInfo().platform.eProductFamily];
1627
1628 TargetDevice targetDevice = {};
1629
1630 targetDevice.coreFamily = pDevice->getHardwareInfo().platform.eRenderCoreFamily;
1631 targetDevice.stepping = pDevice->getHardwareInfo().platform.usRevId;
1632 targetDevice.maxPointerSizeInBytes = sizeof(uintptr_t);
1633
1634 std::string decodeErrors;
1635 std::string decodeWarnings;
1636 auto singleDeviceBinary = unpackSingleDeviceBinary(archive, ConstStringRef(productAbbreviation, strlen(productAbbreviation)), targetDevice,
1637 decodeErrors, decodeWarnings);
1638
1639 EXPECT_FALSE(singleDeviceBinary.debugData.empty());
1640 }
1641
TEST_F(ProgramTests,WhenProgramIsCreatedThenCorrectOclVersionIsInOptions)1642 TEST_F(ProgramTests, WhenProgramIsCreatedThenCorrectOclVersionIsInOptions) {
1643 DebugManagerStateRestore restorer;
1644 DebugManager.flags.DisableStatelessToStatefulOptimization.set(false);
1645
1646 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
1647 auto internalOptions = program.getInitInternalOptions();
1648 if (pClDevice->getEnabledClVersion() == 30) {
1649 EXPECT_TRUE(CompilerOptions::contains(internalOptions, "-ocl-version=300")) << internalOptions;
1650 } else if (pClDevice->getEnabledClVersion() == 21) {
1651 EXPECT_TRUE(CompilerOptions::contains(internalOptions, "-ocl-version=210")) << internalOptions;
1652 } else {
1653 EXPECT_TRUE(CompilerOptions::contains(internalOptions, "-ocl-version=120")) << internalOptions;
1654 }
1655 }
1656
TEST_F(ProgramTests,GivenForcedClVersionWhenProgramIsCreatedThenCorrectOclOptionIsPresent)1657 TEST_F(ProgramTests, GivenForcedClVersionWhenProgramIsCreatedThenCorrectOclOptionIsPresent) {
1658 std::pair<unsigned int, std::string> testedValues[] = {
1659 {0, "-ocl-version=120"},
1660 {12, "-ocl-version=120"},
1661 {21, "-ocl-version=210"},
1662 {30, "-ocl-version=300"}};
1663
1664 for (auto &testedValue : testedValues) {
1665 pClDevice->enabledClVersion = testedValue.first;
1666 MockProgram program{pContext, false, toClDeviceVector(*pClDevice)};
1667 auto internalOptions = program.getInitInternalOptions();
1668 EXPECT_TRUE(CompilerOptions::contains(internalOptions, testedValue.second));
1669 }
1670 }
1671
TEST_F(ProgramTests,GivenStatelessToStatefulIsDisabledWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredOptionIsSet)1672 TEST_F(ProgramTests, GivenStatelessToStatefulIsDisabledWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredOptionIsSet) {
1673 DebugManagerStateRestore restorer;
1674 DebugManager.flags.DisableStatelessToStatefulOptimization.set(true);
1675
1676 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
1677 auto internalOptions = program.getInitInternalOptions();
1678 EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired));
1679 }
1680
TEST_F(ProgramTests,WhenCreatingProgramThenBindlessIsEnabledOnlyIfDebugFlagIsEnabled)1681 TEST_F(ProgramTests, WhenCreatingProgramThenBindlessIsEnabledOnlyIfDebugFlagIsEnabled) {
1682 using namespace testing;
1683 DebugManagerStateRestore restorer;
1684
1685 {
1686
1687 DebugManager.flags.UseBindlessMode.set(0);
1688 MockProgram programNoBindless(pContext, false, toClDeviceVector(*pClDevice));
1689 auto internalOptionsNoBindless = programNoBindless.getInitInternalOptions();
1690 EXPECT_FALSE(CompilerOptions::contains(internalOptionsNoBindless, CompilerOptions::bindlessMode)) << internalOptionsNoBindless;
1691 }
1692 {
1693
1694 DebugManager.flags.UseBindlessMode.set(1);
1695 MockProgram programBindless(pContext, false, toClDeviceVector(*pClDevice));
1696 auto internalOptionsBindless = programBindless.getInitInternalOptions();
1697 EXPECT_TRUE(CompilerOptions::contains(internalOptionsBindless, CompilerOptions::bindlessMode)) << internalOptionsBindless;
1698 }
1699 }
1700
TEST_F(ProgramTests,givenDeviceThatSupportsSharedSystemMemoryAllocationWhenProgramIsCompiledThenItForcesStatelessCompilation)1701 TEST_F(ProgramTests, givenDeviceThatSupportsSharedSystemMemoryAllocationWhenProgramIsCompiledThenItForcesStatelessCompilation) {
1702 pClDevice->deviceInfo.sharedSystemMemCapabilities = CL_UNIFIED_SHARED_MEMORY_ACCESS_INTEL | CL_UNIFIED_SHARED_MEMORY_ATOMIC_ACCESS_INTEL | CL_UNIFIED_SHARED_MEMORY_CONCURRENT_ACCESS_INTEL | CL_UNIFIED_SHARED_MEMORY_CONCURRENT_ATOMIC_ACCESS_INTEL;
1703 pClDevice->getRootDeviceEnvironment().getMutableHardwareInfo()->capabilityTable.sharedSystemMemCapabilities = 1;
1704 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
1705 auto internalOptions = program.getInitInternalOptions();
1706 EXPECT_TRUE(CompilerOptions::contains(internalOptions.c_str(), CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1707 }
1708
TEST_F(ProgramTests,GivenForce32BitAddressessWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet)1709 TEST_F(ProgramTests, GivenForce32BitAddressessWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet) {
1710 cl_int retVal = CL_DEVICE_NOT_FOUND;
1711 auto defaultSetting = DebugManager.flags.DisableStatelessToStatefulOptimization.get();
1712
1713 DebugManager.flags.DisableStatelessToStatefulOptimization.set(false);
1714 if (pDevice) {
1715 const_cast<DeviceInfo *>(&pDevice->getDeviceInfo())->force32BitAddressess = true;
1716 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
1717 auto internalOptions = program.getInitInternalOptions();
1718 if (pDevice->areSharedSystemAllocationsAllowed()) {
1719 EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1720 } else {
1721 EXPECT_FALSE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1722 }
1723 } else {
1724 EXPECT_NE(CL_DEVICE_NOT_FOUND, retVal);
1725 }
1726 DebugManager.flags.DisableStatelessToStatefulOptimization.set(defaultSetting);
1727 }
1728
TEST_F(ProgramTests,Given32bitSupportWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet)1729 TEST_F(ProgramTests, Given32bitSupportWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet) {
1730 auto defaultSetting = DebugManager.flags.DisableStatelessToStatefulOptimization.get();
1731
1732 DebugManager.flags.DisableStatelessToStatefulOptimization.set(false);
1733 std::unique_ptr<MockProgram> program{Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr)};
1734 auto internalOptions = program->getInitInternalOptions();
1735 if ((false == pDevice->areSharedSystemAllocationsAllowed()) && (false == is32bit)) {
1736 EXPECT_FALSE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1737 } else {
1738 EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1739 }
1740 DebugManager.flags.DisableStatelessToStatefulOptimization.set(defaultSetting);
1741 }
1742
TEST_F(ProgramTests,GivenStatelessToStatefulIsDisabledWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet)1743 TEST_F(ProgramTests, GivenStatelessToStatefulIsDisabledWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet) {
1744 auto defaultSetting = DebugManager.flags.DisableStatelessToStatefulOptimization.get();
1745
1746 DebugManager.flags.DisableStatelessToStatefulOptimization.set(true);
1747 std::unique_ptr<MockProgram> program{Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr)};
1748 auto internalOptions = program->getInitInternalOptions();
1749 EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1750 DebugManager.flags.DisableStatelessToStatefulOptimization.set(defaultSetting);
1751 }
1752
TEST_F(ProgramTests,givenProgramWhenItIsCompiledThenItAlwaysHavePreserveVec3TypeInternalOptionSet)1753 TEST_F(ProgramTests, givenProgramWhenItIsCompiledThenItAlwaysHavePreserveVec3TypeInternalOptionSet) {
1754 std::unique_ptr<MockProgram> program(Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr));
1755 auto internalOptions = program->getInitInternalOptions();
1756 EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::preserveVec3Type)) << internalOptions;
1757 }
1758
TEST_F(ProgramTests,Force32BitAddressessWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet)1759 TEST_F(ProgramTests, Force32BitAddressessWhenProgramIsCreatedThenGreaterThan4gbBuffersRequiredIsCorrectlySet) {
1760 auto defaultSetting = DebugManager.flags.DisableStatelessToStatefulOptimization.get();
1761
1762 DebugManager.flags.DisableStatelessToStatefulOptimization.set(false);
1763 const_cast<DeviceInfo *>(&pDevice->getDeviceInfo())->force32BitAddressess = true;
1764 std::unique_ptr<MockProgram> program{Program::createBuiltInFromSource<MockProgram>("", pContext, pContext->getDevices(), nullptr)};
1765 auto internalOptions = program->getInitInternalOptions();
1766 if (is32bit) {
1767 EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1768 } else {
1769 if (false == pDevice->areSharedSystemAllocationsAllowed()) {
1770 EXPECT_FALSE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1771 } else {
1772 EXPECT_TRUE(CompilerOptions::contains(internalOptions, NEO::CompilerOptions::greaterThan4gbBuffersRequired)) << internalOptions;
1773 }
1774 }
1775 DebugManager.flags.DisableStatelessToStatefulOptimization.set(defaultSetting);
1776 }
1777
TEST_F(ProgramTests,GivenStatelessToStatefulBufferOffsetOptimizationWhenProgramIsCreatedThenBufferOffsetArgIsSet)1778 TEST_F(ProgramTests, GivenStatelessToStatefulBufferOffsetOptimizationWhenProgramIsCreatedThenBufferOffsetArgIsSet) {
1779 DebugManagerStateRestore dbgRestorer;
1780 DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(1);
1781 cl_int errorCode = CL_SUCCESS;
1782 const char programSource[] = "program";
1783 const char *programPointer = programSource;
1784 const char **programSources = reinterpret_cast<const char **>(&programPointer);
1785 size_t length = sizeof(programSource);
1786 std::unique_ptr<MockProgram> program(Program::create<MockProgram>(pContext, 1u, programSources, &length, errorCode));
1787 auto internalOptions = program->getInitInternalOptions();
1788
1789 EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg)) << internalOptions;
1790 }
1791
TEST_F(ProgramTests,givenStatelessToStatefullOptimizationOffWHenProgramIsCreatedThenOptimizationStringIsNotPresent)1792 TEST_F(ProgramTests, givenStatelessToStatefullOptimizationOffWHenProgramIsCreatedThenOptimizationStringIsNotPresent) {
1793 DebugManagerStateRestore dbgRestorer;
1794 DebugManager.flags.EnableStatelessToStatefulBufferOffsetOpt.set(0);
1795 cl_int errorCode = CL_SUCCESS;
1796 const char programSource[] = "program";
1797 const char *programPointer = programSource;
1798 const char **programSources = reinterpret_cast<const char **>(&programPointer);
1799 size_t length = sizeof(programSource);
1800 std::unique_ptr<MockProgram> program(Program::create<MockProgram>(pContext, 1u, programSources, &length, errorCode));
1801 auto internalOptions = program->getInitInternalOptions();
1802 EXPECT_FALSE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg)) << internalOptions;
1803 }
1804
TEST_F(ProgramTests,GivenContextWhenCreateProgramThenIncrementContextRefCount)1805 TEST_F(ProgramTests, GivenContextWhenCreateProgramThenIncrementContextRefCount) {
1806 auto initialApiRefCount = pContext->getReference();
1807 auto initialInternalRefCount = pContext->getRefInternalCount();
1808
1809 MockProgram *program = new MockProgram(pContext, false, pContext->getDevices());
1810
1811 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1812 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount + 1);
1813 program->release();
1814 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1815 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount);
1816 }
1817
TEST_F(ProgramTests,GivenContextWhenCreateProgramFromSourceThenIncrementContextRefCount)1818 TEST_F(ProgramTests, GivenContextWhenCreateProgramFromSourceThenIncrementContextRefCount) {
1819 auto initialApiRefCount = pContext->getReference();
1820 auto initialInternalRefCount = pContext->getRefInternalCount();
1821
1822 auto tempProgram = new Program(nullptr, false, pContext->getDevices());
1823 EXPECT_FALSE(tempProgram->getIsBuiltIn());
1824 auto program = new Program(pContext, false, pContext->getDevices());
1825 EXPECT_FALSE(program->getIsBuiltIn());
1826
1827 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1828 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount + 1);
1829 program->release();
1830 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1831 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount);
1832 tempProgram->release();
1833 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1834 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount);
1835 }
1836
TEST_F(ProgramTests,GivenContextWhenCreateBuiltInProgramFromSourceThenDontIncrementContextRefCount)1837 TEST_F(ProgramTests, GivenContextWhenCreateBuiltInProgramFromSourceThenDontIncrementContextRefCount) {
1838 auto initialApiRefCount = pContext->getReference();
1839 auto initialInternalRefCount = pContext->getRefInternalCount();
1840
1841 auto tempProgram = new Program(nullptr, true, pContext->getDevices());
1842 EXPECT_TRUE(tempProgram->getIsBuiltIn());
1843 auto program = new Program(pContext, true, pContext->getDevices());
1844 EXPECT_TRUE(program->getIsBuiltIn());
1845
1846 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1847 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount);
1848 program->release();
1849 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1850 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount);
1851 tempProgram->release();
1852 EXPECT_EQ(pContext->getReference(), initialApiRefCount);
1853 EXPECT_EQ(pContext->getRefInternalCount(), initialInternalRefCount);
1854 }
1855
TEST_F(ProgramTests,WhenBuildingProgramThenPointerToProgramIsReturned)1856 TEST_F(ProgramTests, WhenBuildingProgramThenPointerToProgramIsReturned) {
1857 cl_int retVal = CL_DEVICE_NOT_FOUND;
1858 Program *pProgram = Program::createBuiltInFromSource("", pContext, pContext->getDevices(), &retVal);
1859 EXPECT_NE(nullptr, pProgram);
1860 EXPECT_EQ(CL_SUCCESS, retVal);
1861 delete pProgram;
1862
1863 pProgram = Program::createBuiltInFromSource("", pContext, pContext->getDevices(), nullptr);
1864 EXPECT_NE(nullptr, pProgram);
1865 delete pProgram;
1866 }
1867
TEST_F(ProgramTests,GivenNullBinaryWhenCreatingProgramFromGenBinaryThenInvalidValueErrorIsReturned)1868 TEST_F(ProgramTests, GivenNullBinaryWhenCreatingProgramFromGenBinaryThenInvalidValueErrorIsReturned) {
1869 cl_int retVal = CL_SUCCESS;
1870 Program *pProgram = Program::createBuiltInFromGenBinary(pContext, pContext->getDevices(), nullptr, 0, &retVal);
1871 EXPECT_EQ(nullptr, pProgram);
1872 EXPECT_NE(CL_SUCCESS, retVal);
1873 }
1874
TEST_F(ProgramTests,WhenCreatingProgramFromGenBinaryThenSuccessIsReturned)1875 TEST_F(ProgramTests, WhenCreatingProgramFromGenBinaryThenSuccessIsReturned) {
1876 cl_int retVal = CL_INVALID_BINARY;
1877 char binary[10] = {1, 2, 3, 4, 5, 6, 7, 8, 9, '\0'};
1878 size_t size = 10;
1879
1880 Program *pProgram = Program::createBuiltInFromGenBinary(pContext, pContext->getDevices(), binary, size, &retVal);
1881 EXPECT_NE(nullptr, pProgram);
1882 EXPECT_EQ(CL_SUCCESS, retVal);
1883
1884 EXPECT_EQ((uint32_t)CL_PROGRAM_BINARY_TYPE_EXECUTABLE, (uint32_t)pProgram->getProgramBinaryType(pClDevice));
1885 EXPECT_TRUE(pProgram->getIsBuiltIn());
1886
1887 cl_device_id deviceId = pContext->getDevice(0);
1888 cl_build_status status = 0;
1889 pProgram->getBuildInfo(deviceId, CL_PROGRAM_BUILD_STATUS,
1890 sizeof(cl_build_status), &status, nullptr);
1891 EXPECT_EQ(CL_BUILD_SUCCESS, status);
1892
1893 delete pProgram;
1894 }
1895
TEST_F(ProgramTests,GivenRetValNullPointerWhenCreatingProgramFromGenBinaryThenSuccessIsReturned)1896 TEST_F(ProgramTests, GivenRetValNullPointerWhenCreatingProgramFromGenBinaryThenSuccessIsReturned) {
1897 char binary[10] = {1, 2, 3, 4, 5, 6, 7, 8, 9, '\0'};
1898 size_t size = 10;
1899
1900 Program *pProgram = Program::createBuiltInFromGenBinary(pContext, pContext->getDevices(), binary, size, nullptr);
1901 EXPECT_NE(nullptr, pProgram);
1902 EXPECT_EQ((uint32_t)CL_PROGRAM_BINARY_TYPE_EXECUTABLE, (uint32_t)pProgram->getProgramBinaryType(pClDevice));
1903
1904 cl_device_id deviceId = pContext->getDevice(0);
1905 cl_build_status status = 0;
1906 pProgram->getBuildInfo(deviceId, CL_PROGRAM_BUILD_STATUS,
1907 sizeof(cl_build_status), &status, nullptr);
1908 EXPECT_EQ(CL_BUILD_SUCCESS, status);
1909
1910 delete pProgram;
1911 }
1912
TEST_F(ProgramTests,GivenNullContextWhenCreatingProgramFromGenBinaryThenSuccessIsReturned)1913 TEST_F(ProgramTests, GivenNullContextWhenCreatingProgramFromGenBinaryThenSuccessIsReturned) {
1914 cl_int retVal = CL_INVALID_BINARY;
1915 char binary[10] = {1, 2, 3, 4, 5, 6, 7, 8, 9, '\0'};
1916 size_t size = 10;
1917
1918 Program *pProgram = Program::createBuiltInFromGenBinary(nullptr, toClDeviceVector(*pClDevice), binary, size, &retVal);
1919 EXPECT_NE(nullptr, pProgram);
1920 EXPECT_EQ(CL_SUCCESS, retVal);
1921 EXPECT_EQ((uint32_t)CL_PROGRAM_BINARY_TYPE_EXECUTABLE, (uint32_t)pProgram->getProgramBinaryType(pClDevice));
1922
1923 cl_build_status status = 0;
1924 pProgram->getBuildInfo(pClDevice, CL_PROGRAM_BUILD_STATUS,
1925 sizeof(cl_build_status), &status, nullptr);
1926 EXPECT_EQ(CL_BUILD_SUCCESS, status);
1927
1928 delete pProgram;
1929 }
1930
TEST_F(ProgramTests,givenValidZebinPrepareLinkerInput)1931 TEST_F(ProgramTests, givenValidZebinPrepareLinkerInput) {
1932 ZebinTestData::ValidEmptyProgram zebin;
1933 const std::string validZeInfo = std::string("version :\'") + toString(zeInfoDecoderVersion) + R"===('
1934 kernels:
1935 - name : some_kernel
1936 execution_env :
1937 simd_size : 8
1938 )===";
1939 auto device = std::make_unique<MockClDevice>(MockDevice::createWithNewExecutionEnvironment<MockDevice>(nullptr, mockRootDeviceIndex));
1940 {
1941 auto program = std::make_unique<MockProgram>(nullptr, false, toClDeviceVector(*pClDevice));
1942 program->buildInfos[rootDeviceIndex].unpackedDeviceBinary = makeCopy(zebin.storage.data(), zebin.storage.size());
1943 program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize = zebin.storage.size();
1944
1945 auto retVal = program->processGenBinary(*pClDevice);
1946 EXPECT_EQ(CL_SUCCESS, retVal);
1947 EXPECT_NE(nullptr, program->buildInfos[rootDeviceIndex].linkerInput.get());
1948 }
1949 {
1950 zebin.removeSection(NEO::Elf::SHT_ZEBIN::SHT_ZEBIN_ZEINFO, NEO::Elf::SectionsNamesZebin::zeInfo);
1951 zebin.appendSection(NEO::Elf::SHT_ZEBIN::SHT_ZEBIN_ZEINFO, NEO::Elf::SectionsNamesZebin::zeInfo, ArrayRef<const uint8_t>::fromAny(validZeInfo.data(), validZeInfo.size()));
1952 zebin.appendSection(NEO::Elf::SHT_PROGBITS, NEO::Elf::SectionsNamesZebin::textPrefix.str() + "some_kernel", {});
1953
1954 auto program = std::make_unique<MockProgram>(nullptr, false, toClDeviceVector(*pClDevice));
1955 program->buildInfos[rootDeviceIndex].unpackedDeviceBinary = makeCopy(zebin.storage.data(), zebin.storage.size());
1956 program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize = zebin.storage.size();
1957
1958 auto retVal = program->processGenBinary(*pClDevice);
1959 EXPECT_EQ(CL_SUCCESS, retVal);
1960 EXPECT_NE(nullptr, program->buildInfos[rootDeviceIndex].linkerInput.get());
1961 }
1962 }
1963
TEST_F(ProgramTests,givenProgramFromGenBinaryWhenSLMSizeIsBiggerThenDeviceLimitThenReturnError)1964 TEST_F(ProgramTests, givenProgramFromGenBinaryWhenSLMSizeIsBiggerThenDeviceLimitThenReturnError) {
1965 PatchTokensTestData::ValidProgramWithKernelUsingSlm patchtokensProgram;
1966 patchtokensProgram.slmMutable->TotalInlineLocalMemorySize = static_cast<uint32_t>(pDevice->getDeviceInfo().localMemSize * 2);
1967 patchtokensProgram.recalcTokPtr();
1968 auto program = std::make_unique<MockProgram>(nullptr, false, toClDeviceVector(*pClDevice));
1969 program->buildInfos[rootDeviceIndex].unpackedDeviceBinary = makeCopy(patchtokensProgram.storage.data(), patchtokensProgram.storage.size());
1970 program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize = patchtokensProgram.storage.size();
1971 auto retVal = program->processGenBinary(*pClDevice);
1972
1973 EXPECT_EQ(CL_OUT_OF_RESOURCES, retVal);
1974 }
1975
TEST_F(ProgramTests,givenExistingConstantSurfacesWhenProcessGenBinaryThenCleanupTheSurfaceOnlyForSpecificDevice)1976 TEST_F(ProgramTests, givenExistingConstantSurfacesWhenProcessGenBinaryThenCleanupTheSurfaceOnlyForSpecificDevice) {
1977 PatchTokensTestData::ValidProgramWithKernelUsingSlm patchtokensProgram;
1978
1979 auto program = std::make_unique<MockProgram>(nullptr, false, toClDeviceVector(*pClDevice));
1980
1981 program->buildInfos.resize(2);
1982 program->buildInfos[0].constantSurface = pDevice->getMemoryManager()->allocateGraphicsMemoryWithProperties({rootDeviceIndex, MemoryConstants::cacheLineSize,
1983 GraphicsAllocation::AllocationType::CONSTANT_SURFACE, pDevice->getDeviceBitfield()});
1984 program->buildInfos[1].constantSurface = pDevice->getMemoryManager()->allocateGraphicsMemoryWithProperties({rootDeviceIndex, MemoryConstants::cacheLineSize,
1985 GraphicsAllocation::AllocationType::CONSTANT_SURFACE, pDevice->getDeviceBitfield()});
1986 program->buildInfos[rootDeviceIndex].unpackedDeviceBinary = makeCopy(patchtokensProgram.storage.data(), patchtokensProgram.storage.size());
1987 program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize = patchtokensProgram.storage.size();
1988
1989 auto constantSurface0 = program->buildInfos[0].constantSurface;
1990 EXPECT_NE(nullptr, constantSurface0);
1991 auto constantSurface1 = program->buildInfos[1].constantSurface;
1992 EXPECT_NE(nullptr, constantSurface1);
1993
1994 auto retVal = program->processGenBinary(*pClDevice);
1995
1996 EXPECT_EQ(nullptr, program->buildInfos[0].constantSurface);
1997 EXPECT_EQ(constantSurface1, program->buildInfos[1].constantSurface);
1998
1999 EXPECT_EQ(CL_SUCCESS, retVal);
2000 }
2001
TEST_F(ProgramTests,givenExistingGlobalSurfacesWhenProcessGenBinaryThenCleanupTheSurfaceOnlyForSpecificDevice)2002 TEST_F(ProgramTests, givenExistingGlobalSurfacesWhenProcessGenBinaryThenCleanupTheSurfaceOnlyForSpecificDevice) {
2003 PatchTokensTestData::ValidProgramWithKernelUsingSlm patchtokensProgram;
2004
2005 auto program = std::make_unique<MockProgram>(nullptr, false, toClDeviceVector(*pClDevice));
2006
2007 program->buildInfos.resize(2);
2008 program->buildInfos[0].globalSurface = pDevice->getMemoryManager()->allocateGraphicsMemoryWithProperties({rootDeviceIndex, MemoryConstants::cacheLineSize,
2009 GraphicsAllocation::AllocationType::GLOBAL_SURFACE, pDevice->getDeviceBitfield()});
2010 program->buildInfos[1].globalSurface = pDevice->getMemoryManager()->allocateGraphicsMemoryWithProperties({rootDeviceIndex, MemoryConstants::cacheLineSize,
2011 GraphicsAllocation::AllocationType::GLOBAL_SURFACE, pDevice->getDeviceBitfield()});
2012 program->buildInfos[rootDeviceIndex].unpackedDeviceBinary = makeCopy(patchtokensProgram.storage.data(), patchtokensProgram.storage.size());
2013 program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize = patchtokensProgram.storage.size();
2014
2015 auto globalSurface0 = program->buildInfos[0].globalSurface;
2016 EXPECT_NE(nullptr, globalSurface0);
2017 auto globalSurface1 = program->buildInfos[1].globalSurface;
2018 EXPECT_NE(nullptr, globalSurface1);
2019
2020 auto retVal = program->processGenBinary(*pClDevice);
2021
2022 EXPECT_EQ(nullptr, program->buildInfos[0].globalSurface);
2023 EXPECT_EQ(globalSurface1, program->buildInfos[1].globalSurface);
2024
2025 EXPECT_EQ(CL_SUCCESS, retVal);
2026 }
2027
TEST_F(ProgramTests,GivenNoCompilerInterfaceRootDeviceEnvironmentWhenRebuildingBinaryThenOutOfHostMemoryErrorIsReturned)2028 TEST_F(ProgramTests, GivenNoCompilerInterfaceRootDeviceEnvironmentWhenRebuildingBinaryThenOutOfHostMemoryErrorIsReturned) {
2029 auto pDevice = pContext->getDevice(0);
2030 auto executionEnvironment = pDevice->getExecutionEnvironment();
2031 std::unique_ptr<RootDeviceEnvironment> rootDeviceEnvironment = std::make_unique<NoCompilerInterfaceRootDeviceEnvironment>(*executionEnvironment);
2032 rootDeviceEnvironment->setHwInfo(&pDevice->getHardwareInfo());
2033 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]);
2034 auto program = std::make_unique<MockProgram>(toClDeviceVector(*pDevice));
2035 EXPECT_NE(nullptr, program);
2036
2037 // Load a binary program file
2038 std::string filePath;
2039 retrieveBinaryKernelFilename(filePath, "CopyBuffer_simd16_", ".bin");
2040 size_t binarySize = 0;
2041 auto pBinary = loadDataFromFile(filePath.c_str(), binarySize);
2042 EXPECT_NE(0u, binarySize);
2043
2044 // Create program from loaded binary
2045 cl_int retVal = program->createProgramFromBinary(pBinary.get(), binarySize, *pClDevice);
2046 EXPECT_EQ(CL_SUCCESS, retVal);
2047
2048 // Ask to rebuild program from its IR binary - it should fail (no Compiler Interface)
2049 retVal = program->rebuildProgramFromIr();
2050 EXPECT_EQ(CL_OUT_OF_HOST_MEMORY, retVal);
2051 std::swap(rootDeviceEnvironment, executionEnvironment->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]);
2052 }
2053
TEST_F(ProgramTests,GivenGtpinReraFlagWhenBuildingProgramThenCorrectOptionsAreSet)2054 TEST_F(ProgramTests, GivenGtpinReraFlagWhenBuildingProgramThenCorrectOptionsAreSet) {
2055 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
2056 auto pDevice = pContext->getDevice(0);
2057 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
2058 auto program = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pDevice));
2059 program->sourceCode = "__kernel mock() {}";
2060 program->createdFrom = Program::CreatedFrom::SOURCE;
2061
2062 // Ask to build created program without NEO::CompilerOptions::gtpinRera flag.
2063 cl_int retVal = program->build(program->getDevices(), CompilerOptions::fastRelaxedMath.data(), false);
2064 EXPECT_EQ(CL_SUCCESS, retVal);
2065
2066 // Check build options that were applied
2067 EXPECT_TRUE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::fastRelaxedMath)) << cip->buildOptions;
2068 EXPECT_FALSE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::gtpinRera)) << cip->buildInternalOptions;
2069
2070 // Ask to build created program with NEO::CompilerOptions::gtpinRera flag.
2071 cip->buildOptions.clear();
2072 cip->buildInternalOptions.clear();
2073 retVal = program->build(program->getDevices(), CompilerOptions::concatenate(CompilerOptions::gtpinRera, CompilerOptions::finiteMathOnly).c_str(), false);
2074 EXPECT_EQ(CL_SUCCESS, retVal);
2075
2076 // Check build options that were applied
2077 EXPECT_FALSE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::fastRelaxedMath)) << cip->buildOptions;
2078 EXPECT_TRUE(CompilerOptions::contains(cip->buildOptions, CompilerOptions::finiteMathOnly)) << cip->buildOptions;
2079 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, CompilerOptions::gtpinRera)) << cip->buildInternalOptions;
2080 }
2081
TEST_F(ProgramTests,GivenFailingGenBinaryProgramWhenRebuildingBinaryThenInvalidBinaryErrorIsReturned)2082 TEST_F(ProgramTests, GivenFailingGenBinaryProgramWhenRebuildingBinaryThenInvalidBinaryErrorIsReturned) {
2083
2084 cl_int retVal;
2085
2086 auto program = std::make_unique<FailingGenBinaryProgram>(toClDeviceVector(*pClDevice));
2087 EXPECT_NE(nullptr, program);
2088
2089 // Load a binary program file
2090 std::string filePath;
2091 retrieveBinaryKernelFilename(filePath, "CopyBuffer_simd16_", ".bin");
2092 size_t binarySize = 0;
2093 auto pBinary = loadDataFromFile(filePath.c_str(), binarySize);
2094 EXPECT_NE(0u, binarySize);
2095
2096 // Create program from loaded binary
2097 retVal = program->createProgramFromBinary(pBinary.get(), binarySize, *pClDevice);
2098 EXPECT_EQ(CL_SUCCESS, retVal);
2099
2100 // Ask to rebuild program from its IR binary - it should fail (simulated invalid binary)
2101 retVal = program->rebuildProgramFromIr();
2102 EXPECT_EQ(CL_INVALID_BINARY, retVal);
2103 }
2104
TEST_F(ProgramTests,GivenZeroPrivateSizeInBlockWhenAllocateBlockProvateSurfacesCalledThenNoSurfaceIsCreated)2105 TEST_F(ProgramTests, GivenZeroPrivateSizeInBlockWhenAllocateBlockProvateSurfacesCalledThenNoSurfaceIsCreated) {
2106 MockProgram *program = new MockProgram(pContext, false, toClDeviceVector(*pClDevice));
2107
2108 uint32_t crossThreadOffsetBlock = 0;
2109
2110 auto infoBlock = new MockKernelInfo;
2111 infoBlock->setPrivateMemory(0, false, 8, crossThreadOffsetBlock, 0);
2112
2113 program->blockKernelManager->addBlockKernelInfo(infoBlock);
2114
2115 program->allocateBlockPrivateSurfaces(*pClDevice);
2116
2117 EXPECT_EQ(nullptr, program->getBlockKernelManager()->getPrivateSurface(0));
2118
2119 delete program;
2120 }
2121
TEST_F(ProgramTests,GivenNonZeroPrivateSizeInBlockWhenAllocateBlockProvateSurfacesCalledThenSurfaceIsCreated)2122 TEST_F(ProgramTests, GivenNonZeroPrivateSizeInBlockWhenAllocateBlockProvateSurfacesCalledThenSurfaceIsCreated) {
2123 MockProgram *program = new MockProgram(pContext, false, toClDeviceVector(*pClDevice));
2124
2125 uint32_t crossThreadOffsetBlock = 0;
2126
2127 auto infoBlock = new MockKernelInfo;
2128 infoBlock->setPrivateMemory(1000, false, 8, crossThreadOffsetBlock, 0);
2129
2130 program->blockKernelManager->addBlockKernelInfo(infoBlock);
2131
2132 program->allocateBlockPrivateSurfaces(*pClDevice);
2133
2134 EXPECT_NE(nullptr, program->getBlockKernelManager()->getPrivateSurface(0));
2135
2136 delete program;
2137 }
2138
TEST_F(ProgramTests,GivenNonZeroPrivateSizeInBlockWhenAllocateBlockProvateSurfacesCalledThenSecondSurfaceIsNotCreated)2139 TEST_F(ProgramTests, GivenNonZeroPrivateSizeInBlockWhenAllocateBlockProvateSurfacesCalledThenSecondSurfaceIsNotCreated) {
2140 MockProgram *program = new MockProgram(pContext, false, toClDeviceVector(*pClDevice));
2141
2142 uint32_t crossThreadOffsetBlock = 0;
2143
2144 auto infoBlock = new MockKernelInfo;
2145 infoBlock->setPrivateMemory(1000, false, 8, crossThreadOffsetBlock, 0);
2146
2147 program->blockKernelManager->addBlockKernelInfo(infoBlock);
2148
2149 program->allocateBlockPrivateSurfaces(*pClDevice);
2150
2151 GraphicsAllocation *privateSurface = program->getBlockKernelManager()->getPrivateSurface(0);
2152
2153 EXPECT_NE(nullptr, privateSurface);
2154
2155 program->allocateBlockPrivateSurfaces(*pClDevice);
2156
2157 GraphicsAllocation *privateSurface2 = program->getBlockKernelManager()->getPrivateSurface(0);
2158
2159 EXPECT_EQ(privateSurface, privateSurface2);
2160
2161 delete program;
2162 }
2163
TEST_F(ProgramTests,givenProgramWithBlockKernelsWhenfreeBlockResourcesisCalledThenFreeGraphhicsAllocationsFromBlockKernelManagerIsCalled)2164 TEST_F(ProgramTests, givenProgramWithBlockKernelsWhenfreeBlockResourcesisCalledThenFreeGraphhicsAllocationsFromBlockKernelManagerIsCalled) {
2165 MockProgram *program = new MockProgram(pContext, false, toClDeviceVector(*pClDevice));
2166
2167 uint32_t crossThreadOffsetBlock = 0;
2168
2169 auto infoBlock = new MockKernelInfo;
2170 infoBlock->setPrivateMemory(1000, false, 8, crossThreadOffsetBlock, 0);
2171
2172 program->blockKernelManager->addBlockKernelInfo(infoBlock);
2173
2174 GraphicsAllocation *privateSurface = pDevice->getMemoryManager()->allocateGraphicsMemoryWithProperties(MockAllocationProperties{pDevice->getRootDeviceIndex(), MemoryConstants::pageSize});
2175 EXPECT_NE(nullptr, privateSurface);
2176
2177 program->getBlockKernelManager()->pushPrivateSurface(privateSurface, 0);
2178
2179 program->freeBlockResources();
2180
2181 delete program;
2182 }
2183
2184 class Program32BitTests : public ProgramTests {
2185 public:
SetUp()2186 void SetUp() override {
2187 DebugManager.flags.Force32bitAddressing.set(true);
2188 ProgramTests::SetUp();
2189 }
TearDown()2190 void TearDown() override {
2191 ProgramTests::TearDown();
2192 DebugManager.flags.Force32bitAddressing.set(false);
2193 }
2194 };
2195
TEST_F(Program32BitTests,givenDeviceWithForce32BitAddressingOnWhenBuiltinIsCreatedThenNoFlagsArePassedAsInternalOptions)2196 TEST_F(Program32BitTests, givenDeviceWithForce32BitAddressingOnWhenBuiltinIsCreatedThenNoFlagsArePassedAsInternalOptions) {
2197 MockProgram program(toClDeviceVector(*pClDevice));
2198 auto internalOptions = program.getInitInternalOptions();
2199 EXPECT_THAT(internalOptions, testing::HasSubstr(std::string("")));
2200 }
2201
TEST_F(Program32BitTests,givenDeviceWithForce32BitAddressingOnWhenProgramIsCreatedThen32bitFlagIsPassedAsInternalOption)2202 TEST_F(Program32BitTests, givenDeviceWithForce32BitAddressingOnWhenProgramIsCreatedThen32bitFlagIsPassedAsInternalOption) {
2203 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2204 auto internalOptions = program.getInitInternalOptions();
2205 std::string s1 = internalOptions;
2206 size_t pos = s1.find(NEO::CompilerOptions::arch32bit.data());
2207 if constexpr (is64bit) {
2208 EXPECT_NE(pos, std::string::npos);
2209 } else {
2210 EXPECT_EQ(pos, std::string::npos);
2211 }
2212 }
2213
HWTEST_F(ProgramTests,givenNewProgramThenStatelessToStatefulBufferOffsetOptimizationIsMatchingThePlatformEnablingStatus)2214 HWTEST_F(ProgramTests, givenNewProgramThenStatelessToStatefulBufferOffsetOptimizationIsMatchingThePlatformEnablingStatus) {
2215 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2216 auto internalOptions = program.getInitInternalOptions();
2217
2218 if (HwHelperHw<FamilyType>::get().isStatelesToStatefullWithOffsetSupported()) {
2219 EXPECT_TRUE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg));
2220 } else {
2221 EXPECT_FALSE(CompilerOptions::contains(internalOptions, CompilerOptions::hasBufferOffsetArg));
2222 }
2223 }
2224
TEST(ProgramTest,givenImagesSupportedWhenCreatingProgramThenInternalOptionsAreCorrectlyInitialized)2225 TEST(ProgramTest, givenImagesSupportedWhenCreatingProgramThenInternalOptionsAreCorrectlyInitialized) {
2226 VariableBackup<bool> supportsImagesCapability{&defaultHwInfo->capabilityTable.supportsImages};
2227
2228 for (auto areImagesSupported : ::testing::Bool()) {
2229 supportsImagesCapability = areImagesSupported;
2230 UltClDeviceFactory clDeviceFactory{1, 0};
2231 MockContext context{clDeviceFactory.rootDevices[0]};
2232 MockProgram program(&context, false, toClDeviceVector(*clDeviceFactory.rootDevices[0]));
2233
2234 auto internalOptions = program.getInitInternalOptions();
2235 EXPECT_EQ(areImagesSupported, CompilerOptions::contains(internalOptions, CompilerOptions::enableImageSupport));
2236 }
2237 }
2238
2239 template <int32_t ErrCodeToReturn, bool spirv = true>
2240 struct CreateProgramFromBinaryMock : public MockProgram {
2241 using MockProgram::MockProgram;
2242
createProgramFromBinaryCreateProgramFromBinaryMock2243 cl_int createProgramFromBinary(const void *pBinary,
2244 size_t binarySize, ClDevice &clDevice) override {
2245 this->irBinary.reset(new char[binarySize]);
2246 this->irBinarySize = binarySize;
2247 this->isSpirV = spirv;
2248 memcpy_s(this->irBinary.get(), binarySize, pBinary, binarySize);
2249 return ErrCodeToReturn;
2250 }
2251 };
2252
TEST_F(ProgramTests,GivenFailedBinaryWhenCreatingFromIlThenInvalidBinaryErrorIsReturned)2253 TEST_F(ProgramTests, GivenFailedBinaryWhenCreatingFromIlThenInvalidBinaryErrorIsReturned) {
2254 REQUIRE_OCL_21_OR_SKIP(pContext);
2255 const uint32_t notSpirv[16] = {0xDEADBEEF};
2256 cl_int retVal = CL_SUCCESS;
2257 auto prog = Program::createFromIL<CreateProgramFromBinaryMock<CL_INVALID_BINARY>>(pContext, reinterpret_cast<const void *>(notSpirv), sizeof(notSpirv), retVal);
2258 EXPECT_EQ(nullptr, prog);
2259 EXPECT_EQ(CL_INVALID_BINARY, retVal);
2260 }
2261
TEST_F(ProgramTests,GivenSuccessfullyBuiltBinaryWhenCreatingFromIlThenValidProgramIsReturned)2262 TEST_F(ProgramTests, GivenSuccessfullyBuiltBinaryWhenCreatingFromIlThenValidProgramIsReturned) {
2263 REQUIRE_OCL_21_OR_SKIP(pContext);
2264 const uint32_t spirv[16] = {0x03022307};
2265 cl_int retVal = CL_SUCCESS;
2266 auto prog = Program::createFromIL<CreateProgramFromBinaryMock<CL_SUCCESS>>(pContext, reinterpret_cast<const void *>(spirv), sizeof(spirv), retVal);
2267 ASSERT_NE(nullptr, prog);
2268 EXPECT_EQ(CL_SUCCESS, retVal);
2269 prog->release();
2270 }
2271
TEST_F(ProgramTests,givenProgramCreatedFromILWhenCompileIsCalledThenReuseTheILInsteadOfCallingCompilerInterface)2272 TEST_F(ProgramTests, givenProgramCreatedFromILWhenCompileIsCalledThenReuseTheILInsteadOfCallingCompilerInterface) {
2273 REQUIRE_OCL_21_OR_SKIP(pContext);
2274 const uint32_t spirv[16] = {0x03022307};
2275 cl_int errCode = 0;
2276 auto pProgram = Program::createFromIL<MockProgram>(pContext, reinterpret_cast<const void *>(spirv), sizeof(spirv), errCode);
2277 ASSERT_NE(nullptr, pProgram);
2278 auto debugVars = NEO::getIgcDebugVars();
2279 debugVars.forceBuildFailure = true;
2280 gEnvironment->fclPushDebugVars(debugVars);
2281 auto compilerErr = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
2282 EXPECT_EQ(CL_SUCCESS, compilerErr);
2283 gEnvironment->fclPopDebugVars();
2284 pProgram->release();
2285 }
2286
TEST_F(ProgramTests,givenProgramCreatedFromIntermediateBinaryRepresentationWhenCompileIsCalledThenReuseTheILInsteadOfCallingCompilerInterface)2287 TEST_F(ProgramTests, givenProgramCreatedFromIntermediateBinaryRepresentationWhenCompileIsCalledThenReuseTheILInsteadOfCallingCompilerInterface) {
2288 const uint32_t spirv[16] = {0x03022307};
2289 cl_int errCode = 0;
2290 size_t lengths = sizeof(spirv);
2291 const unsigned char *binaries[1] = {reinterpret_cast<const unsigned char *>(spirv)};
2292 auto pProgram = Program::create<MockProgram>(pContext, pContext->getDevices(), &lengths, binaries, nullptr, errCode);
2293 ASSERT_NE(nullptr, pProgram);
2294 auto debugVars = NEO::getIgcDebugVars();
2295 debugVars.forceBuildFailure = true;
2296 gEnvironment->fclPushDebugVars(debugVars);
2297 auto compilerErr = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
2298 EXPECT_EQ(CL_SUCCESS, compilerErr);
2299 gEnvironment->fclPopDebugVars();
2300 pProgram->release();
2301 }
2302
TEST_F(ProgramTests,GivenIlIsNullptrWhenCreatingFromIlThenInvalidBinaryErrorIsReturned)2303 TEST_F(ProgramTests, GivenIlIsNullptrWhenCreatingFromIlThenInvalidBinaryErrorIsReturned) {
2304 REQUIRE_OCL_21_OR_SKIP(pContext);
2305 cl_int retVal = CL_SUCCESS;
2306 auto prog = Program::createFromIL<CreateProgramFromBinaryMock<CL_INVALID_BINARY>>(pContext, nullptr, 16, retVal);
2307 EXPECT_EQ(nullptr, prog);
2308 EXPECT_EQ(CL_INVALID_BINARY, retVal);
2309 }
2310
TEST_F(ProgramTests,GivenIlSizeZeroWhenCreatingFromIlThenInvalidBinaryErrorIsReturned)2311 TEST_F(ProgramTests, GivenIlSizeZeroWhenCreatingFromIlThenInvalidBinaryErrorIsReturned) {
2312 REQUIRE_OCL_21_OR_SKIP(pContext);
2313 const uint32_t spirv[16] = {0x03022307};
2314 cl_int retVal = CL_SUCCESS;
2315 auto prog = Program::createFromIL<CreateProgramFromBinaryMock<CL_INVALID_BINARY>>(pContext, reinterpret_cast<const void *>(spirv), 0, retVal);
2316 EXPECT_EQ(nullptr, prog);
2317 EXPECT_EQ(CL_INVALID_BINARY, retVal);
2318 }
2319
TEST_F(ProgramTests,WhenCreatingFromIlThenIsSpirvIsSetCorrectly)2320 TEST_F(ProgramTests, WhenCreatingFromIlThenIsSpirvIsSetCorrectly) {
2321 REQUIRE_OCL_21_OR_SKIP(pContext);
2322 const uint32_t spirv[16] = {0x03022307};
2323 cl_int retVal = CL_SUCCESS;
2324 auto prog = Program::createFromIL<Program>(pContext, reinterpret_cast<const void *>(spirv), sizeof(spirv), retVal);
2325 EXPECT_NE(nullptr, prog);
2326 EXPECT_EQ(CL_SUCCESS, retVal);
2327 EXPECT_TRUE(prog->getIsSpirV());
2328 prog->release();
2329
2330 const char llvmBc[16] = {'B', 'C', '\xc0', '\xde'};
2331 prog = Program::createFromIL<Program>(pContext, reinterpret_cast<const void *>(llvmBc), sizeof(llvmBc), retVal);
2332 EXPECT_NE(nullptr, prog);
2333 EXPECT_EQ(CL_SUCCESS, retVal);
2334 EXPECT_FALSE(prog->getIsSpirV());
2335 prog->release();
2336 }
2337
2338 static const uint8_t llvmBinary[] = "BC\xc0\xde ";
2339
TEST(isValidLlvmBinary,whenLlvmMagicWasFoundThenBinaryIsValidLLvm)2340 TEST(isValidLlvmBinary, whenLlvmMagicWasFoundThenBinaryIsValidLLvm) {
2341 EXPECT_TRUE(NEO::isLlvmBitcode(llvmBinary));
2342 }
2343
TEST(isValidLlvmBinary,whenBinaryIsNullptrThenBinaryIsNotValidLLvm)2344 TEST(isValidLlvmBinary, whenBinaryIsNullptrThenBinaryIsNotValidLLvm) {
2345 EXPECT_FALSE(NEO::isLlvmBitcode(ArrayRef<const uint8_t>()));
2346 }
2347
TEST(isValidLlvmBinary,whenBinaryIsShorterThanLllvMagicThenBinaryIsNotValidLLvm)2348 TEST(isValidLlvmBinary, whenBinaryIsShorterThanLllvMagicThenBinaryIsNotValidLLvm) {
2349 EXPECT_FALSE(NEO::isLlvmBitcode(ArrayRef<const uint8_t>(llvmBinary, 2)));
2350 }
2351
TEST(isValidLlvmBinary,whenBinaryDoesNotContainLllvMagicThenBinaryIsNotValidLLvm)2352 TEST(isValidLlvmBinary, whenBinaryDoesNotContainLllvMagicThenBinaryIsNotValidLLvm) {
2353 const uint8_t notLlvmBinary[] = "ABCDEFGHIJKLMNO";
2354 EXPECT_FALSE(NEO::isLlvmBitcode(notLlvmBinary));
2355 }
2356
2357 const uint32_t spirv[16] = {0x03022307};
2358 const uint32_t spirvInvEndianes[16] = {0x07230203};
2359
TEST(isValidSpirvBinary,whenSpirvMagicWasFoundThenBinaryIsValidSpirv)2360 TEST(isValidSpirvBinary, whenSpirvMagicWasFoundThenBinaryIsValidSpirv) {
2361 EXPECT_TRUE(NEO::isSpirVBitcode(ArrayRef<const uint8_t>(reinterpret_cast<const uint8_t *>(&spirv), sizeof(spirv))));
2362 EXPECT_TRUE(NEO::isSpirVBitcode(ArrayRef<const uint8_t>(reinterpret_cast<const uint8_t *>(&spirvInvEndianes), sizeof(spirvInvEndianes))));
2363 }
2364
TEST(isValidSpirvBinary,whenBinaryIsNullptrThenBinaryIsNotValidLLvm)2365 TEST(isValidSpirvBinary, whenBinaryIsNullptrThenBinaryIsNotValidLLvm) {
2366 EXPECT_FALSE(NEO::isSpirVBitcode(ArrayRef<const uint8_t>()));
2367 }
2368
TEST(isValidSpirvBinary,whenBinaryIsShorterThanLllvMagicThenBinaryIsNotValidLLvm)2369 TEST(isValidSpirvBinary, whenBinaryIsShorterThanLllvMagicThenBinaryIsNotValidLLvm) {
2370 EXPECT_FALSE(NEO::isSpirVBitcode(ArrayRef<const uint8_t>(reinterpret_cast<const uint8_t *>(&spirvInvEndianes), 2)));
2371 }
2372
TEST(isValidSpirvBinary,whenBinaryDoesNotContainLllvMagicThenBinaryIsNotValidLLvm)2373 TEST(isValidSpirvBinary, whenBinaryDoesNotContainLllvMagicThenBinaryIsNotValidLLvm) {
2374 const uint8_t notSpirvBinary[] = "ABCDEFGHIJKLMNO";
2375 EXPECT_FALSE(NEO::isSpirVBitcode(notSpirvBinary));
2376 }
2377
TEST_F(ProgramTests,WhenLinkingTwoValidSpirvProgramsThenValidProgramIsReturned)2378 TEST_F(ProgramTests, WhenLinkingTwoValidSpirvProgramsThenValidProgramIsReturned) {
2379 REQUIRE_OCL_21_OR_SKIP(pContext);
2380 const uint32_t spirv[16] = {0x03022307};
2381 cl_int errCode = CL_SUCCESS;
2382
2383 auto node1 = Program::createFromIL<CreateProgramFromBinaryMock<CL_SUCCESS, false>>(pContext, reinterpret_cast<const void *>(spirv), sizeof(spirv), errCode);
2384 ASSERT_NE(nullptr, node1);
2385 EXPECT_EQ(CL_SUCCESS, errCode);
2386
2387 auto node2 = Program::createFromIL<CreateProgramFromBinaryMock<CL_SUCCESS>>(pContext, reinterpret_cast<const void *>(spirv), sizeof(spirv), errCode);
2388 ASSERT_NE(nullptr, node2);
2389 EXPECT_EQ(CL_SUCCESS, errCode);
2390
2391 auto prog = Program::createFromIL<CreateProgramFromBinaryMock<CL_SUCCESS>>(pContext, reinterpret_cast<const void *>(spirv), sizeof(spirv), errCode);
2392 ASSERT_NE(nullptr, prog);
2393 EXPECT_EQ(CL_SUCCESS, errCode);
2394
2395 cl_program linkNodes[] = {node1, node2};
2396 errCode = prog->link(prog->getDevices(), nullptr, 2, linkNodes);
2397 EXPECT_EQ(CL_SUCCESS, errCode);
2398
2399 prog->release();
2400 node2->release();
2401 node1->release();
2402 }
2403
TEST_F(ProgramTests,givenSeparateBlockKernelsWhenNoParentAndSubgroupKernelsThenSeparateNoneKernel)2404 TEST_F(ProgramTests, givenSeparateBlockKernelsWhenNoParentAndSubgroupKernelsThenSeparateNoneKernel) {
2405 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2406
2407 EXPECT_EQ(0u, program.getKernelInfoArray(rootDeviceIndex).size());
2408 EXPECT_EQ(0u, program.getParentKernelInfoArray(rootDeviceIndex).size());
2409 EXPECT_EQ(0u, program.getSubgroupKernelInfoArray(rootDeviceIndex).size());
2410
2411 program.separateBlockKernels(rootDeviceIndex);
2412
2413 EXPECT_EQ(0u, program.getKernelInfoArray(rootDeviceIndex).size());
2414 EXPECT_EQ(0u, program.getBlockKernelManager()->getCount());
2415 }
2416
TEST_F(ProgramTests,givenSeparateBlockKernelsWhenRegularKernelsThenSeparateNoneKernel)2417 TEST_F(ProgramTests, givenSeparateBlockKernelsWhenRegularKernelsThenSeparateNoneKernel) {
2418 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2419
2420 auto pRegularKernel1Info = new KernelInfo();
2421 pRegularKernel1Info->kernelDescriptor.kernelMetadata.kernelName = "regular_kernel_1";
2422 program.getKernelInfoArray(rootDeviceIndex).push_back(pRegularKernel1Info);
2423
2424 auto pRegularKernel2Info = new KernelInfo();
2425 pRegularKernel2Info->kernelDescriptor.kernelMetadata.kernelName = "regular_kernel_2";
2426 program.getKernelInfoArray(rootDeviceIndex).push_back(pRegularKernel2Info);
2427
2428 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2429
2430 program.separateBlockKernels(rootDeviceIndex);
2431
2432 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2433 EXPECT_EQ(0, strcmp("regular_kernel_1", program.getKernelInfoArray(rootDeviceIndex).at(0)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2434 EXPECT_EQ(0, strcmp("regular_kernel_2", program.getKernelInfoArray(rootDeviceIndex).at(1)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2435
2436 EXPECT_EQ(0u, program.getBlockKernelManager()->getCount());
2437 }
2438
TEST_F(ProgramTests,givenSeparateBlockKernelsWhenChildLikeKernelWithoutParentKernelThenSeparateNoneKernel)2439 TEST_F(ProgramTests, givenSeparateBlockKernelsWhenChildLikeKernelWithoutParentKernelThenSeparateNoneKernel) {
2440 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2441
2442 auto pParentKernelInfo = new KernelInfo();
2443 pParentKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "another_parent_kernel";
2444 program.getKernelInfoArray(rootDeviceIndex).push_back(pParentKernelInfo);
2445 program.getParentKernelInfoArray(rootDeviceIndex).push_back(pParentKernelInfo);
2446
2447 auto pChildKernelInfo = new KernelInfo();
2448 pChildKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "childlike_kernel_dispatch_0";
2449 program.getKernelInfoArray(rootDeviceIndex).push_back(pChildKernelInfo);
2450
2451 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2452 EXPECT_EQ(1u, program.getParentKernelInfoArray(rootDeviceIndex).size());
2453
2454 program.separateBlockKernels(rootDeviceIndex);
2455
2456 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2457 EXPECT_EQ(0, strcmp("another_parent_kernel", program.getKernelInfoArray(rootDeviceIndex).at(0)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2458 EXPECT_EQ(0, strcmp("childlike_kernel_dispatch_0", program.getKernelInfoArray(rootDeviceIndex).at(1)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2459
2460 EXPECT_EQ(0u, program.getBlockKernelManager()->getCount());
2461 }
2462
TEST_F(ProgramTests,givenSeparateBlockKernelsWhenChildLikeKernelWithoutSubgroupKernelThenSeparateNoneKernel)2463 TEST_F(ProgramTests, givenSeparateBlockKernelsWhenChildLikeKernelWithoutSubgroupKernelThenSeparateNoneKernel) {
2464 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2465
2466 auto pSubgroupKernelInfo = new KernelInfo();
2467 pSubgroupKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "another_subgroup_kernel";
2468 program.getKernelInfoArray(rootDeviceIndex).push_back(pSubgroupKernelInfo);
2469 program.getSubgroupKernelInfoArray(rootDeviceIndex).push_back(pSubgroupKernelInfo);
2470
2471 auto pChildKernelInfo = new KernelInfo();
2472 pChildKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "childlike_kernel_dispatch_0";
2473 program.getKernelInfoArray(rootDeviceIndex).push_back(pChildKernelInfo);
2474
2475 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2476 EXPECT_EQ(1u, program.getSubgroupKernelInfoArray(rootDeviceIndex).size());
2477
2478 program.separateBlockKernels(rootDeviceIndex);
2479
2480 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2481 EXPECT_EQ(0, strcmp("another_subgroup_kernel", program.getKernelInfoArray(rootDeviceIndex).at(0)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2482 EXPECT_EQ(0, strcmp("childlike_kernel_dispatch_0", program.getKernelInfoArray(rootDeviceIndex).at(1)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2483
2484 EXPECT_EQ(0u, program.getBlockKernelManager()->getCount());
2485 }
2486
TEST_F(ProgramTests,givenSeparateBlockKernelsWhenParentKernelWithChildKernelThenSeparateChildKernel)2487 TEST_F(ProgramTests, givenSeparateBlockKernelsWhenParentKernelWithChildKernelThenSeparateChildKernel) {
2488 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2489
2490 auto pParentKernelInfo = new KernelInfo();
2491 pParentKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "parent_kernel";
2492 program.getKernelInfoArray(rootDeviceIndex).push_back(pParentKernelInfo);
2493 program.getParentKernelInfoArray(rootDeviceIndex).push_back(pParentKernelInfo);
2494
2495 auto pChildKernelInfo = new KernelInfo();
2496 pChildKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "parent_kernel_dispatch_0";
2497 program.getKernelInfoArray(rootDeviceIndex).push_back(pChildKernelInfo);
2498
2499 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2500 EXPECT_EQ(1u, program.getParentKernelInfoArray(rootDeviceIndex).size());
2501
2502 program.separateBlockKernels(rootDeviceIndex);
2503
2504 EXPECT_EQ(1u, program.getKernelInfoArray(rootDeviceIndex).size());
2505 EXPECT_EQ(0, strcmp("parent_kernel", program.getKernelInfoArray(rootDeviceIndex).at(0)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2506
2507 EXPECT_EQ(1u, program.getBlockKernelManager()->getCount());
2508 EXPECT_EQ(0, strcmp("parent_kernel_dispatch_0", program.getBlockKernelManager()->getBlockKernelInfo(0)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2509 }
2510
TEST_F(ProgramTests,givenSeparateBlockKernelsWhenSubgroupKernelWithChildKernelThenSeparateChildKernel)2511 TEST_F(ProgramTests, givenSeparateBlockKernelsWhenSubgroupKernelWithChildKernelThenSeparateChildKernel) {
2512 MockProgram program(pContext, false, toClDeviceVector(*pClDevice));
2513
2514 auto pSubgroupKernelInfo = new KernelInfo();
2515 pSubgroupKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "subgroup_kernel";
2516 program.getKernelInfoArray(rootDeviceIndex).push_back(pSubgroupKernelInfo);
2517 program.getSubgroupKernelInfoArray(rootDeviceIndex).push_back(pSubgroupKernelInfo);
2518
2519 auto pChildKernelInfo = new KernelInfo();
2520 pChildKernelInfo->kernelDescriptor.kernelMetadata.kernelName = "subgroup_kernel_dispatch_0";
2521 program.getKernelInfoArray(rootDeviceIndex).push_back(pChildKernelInfo);
2522
2523 EXPECT_EQ(2u, program.getKernelInfoArray(rootDeviceIndex).size());
2524 EXPECT_EQ(1u, program.getSubgroupKernelInfoArray(rootDeviceIndex).size());
2525
2526 program.separateBlockKernels(rootDeviceIndex);
2527
2528 EXPECT_EQ(1u, program.getKernelInfoArray(rootDeviceIndex).size());
2529 EXPECT_EQ(0, strcmp("subgroup_kernel", program.getKernelInfoArray(rootDeviceIndex).at(0)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2530
2531 EXPECT_EQ(1u, program.getBlockKernelManager()->getCount());
2532 EXPECT_EQ(0, strcmp("subgroup_kernel_dispatch_0", program.getBlockKernelManager()->getBlockKernelInfo(0)->kernelDescriptor.kernelMetadata.kernelName.c_str()));
2533 }
2534
TEST(ProgramDestructionTests,givenProgramUsingDeviceWhenItIsDestroyedAfterPlatfromCleanupThenItIsCleanedUpProperly)2535 TEST(ProgramDestructionTests, givenProgramUsingDeviceWhenItIsDestroyedAfterPlatfromCleanupThenItIsCleanedUpProperly) {
2536 initPlatform();
2537 auto device = platform()->getClDevice(0);
2538 MockContext *context = new MockContext(device, false);
2539 MockProgram *pProgram = new MockProgram(context, false, toClDeviceVector(*device));
2540 auto globalAllocation = device->getMemoryManager()->allocateGraphicsMemoryWithProperties(MockAllocationProperties{device->getRootDeviceIndex(), MemoryConstants::pageSize});
2541 pProgram->setGlobalSurface(globalAllocation);
2542
2543 platformsImpl->clear();
2544 EXPECT_EQ(1, device->getRefInternalCount());
2545 EXPECT_EQ(1, pProgram->getRefInternalCount());
2546 context->decRefInternal();
2547 pProgram->decRefInternal();
2548 }
2549
TEST_F(ProgramTests,givenProgramWithSpirvWhenRebuildProgramIsCalledThenSpirvPathIsTaken)2550 TEST_F(ProgramTests, givenProgramWithSpirvWhenRebuildProgramIsCalledThenSpirvPathIsTaken) {
2551 auto compilerInterface = new MockCompilerInterface();
2552 auto compilerMain = new MockCIFMain();
2553 compilerInterface->setFclMain(compilerMain);
2554 compilerMain->Retain();
2555 compilerInterface->setIgcMain(compilerMain);
2556 compilerMain->setDefaultCreatorFunc<NEO::MockIgcOclDeviceCtx>(NEO::MockIgcOclDeviceCtx::Create);
2557 compilerMain->setDefaultCreatorFunc<NEO::MockFclOclDeviceCtx>(NEO::MockFclOclDeviceCtx::Create);
2558 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(compilerInterface);
2559
2560 std::string receivedInput;
2561 MockCompilerDebugVars debugVars = {};
2562 debugVars.receivedInput = &receivedInput;
2563 debugVars.forceBuildFailure = true;
2564 gEnvironment->igcPushDebugVars(debugVars);
2565 std::unique_ptr<void, void (*)(void *)> igcDebugVarsAutoPop{&gEnvironment, [](void *) { gEnvironment->igcPopDebugVars(); }};
2566
2567 auto program = clUniquePtr(new MockProgram(toClDeviceVector(*pClDevice)));
2568 uint32_t spirv[16] = {0x03022307, 0x23471113, 0x17192329};
2569 program->irBinary = makeCopy(spirv, sizeof(spirv));
2570 program->irBinarySize = sizeof(spirv);
2571 program->isSpirV = true;
2572 auto buildRet = program->rebuildProgramFromIr();
2573 EXPECT_NE(CL_SUCCESS, buildRet);
2574 ASSERT_EQ(sizeof(spirv), receivedInput.size());
2575 EXPECT_EQ(0, memcmp(spirv, receivedInput.c_str(), receivedInput.size()));
2576 ASSERT_EQ(1U, compilerInterface->requestedTranslationCtxs.size());
2577 EXPECT_EQ(IGC::CodeType::spirV, compilerInterface->requestedTranslationCtxs[0].first);
2578 EXPECT_EQ(IGC::CodeType::oclGenBin, compilerInterface->requestedTranslationCtxs[0].second);
2579 }
2580
TEST_F(ProgramTests,givenProgramWithSpirvWhenRebuildIsCalledThenRebuildWarningIsIssued)2581 TEST_F(ProgramTests, givenProgramWithSpirvWhenRebuildIsCalledThenRebuildWarningIsIssued) {
2582 const auto program{clUniquePtr(new MockProgram(toClDeviceVector(*pClDevice)))};
2583 uint32_t spirv[16] = {0x03022307, 0x23471113, 0x17192329};
2584 program->irBinary = makeCopy(spirv, sizeof(spirv));
2585 program->irBinarySize = sizeof(spirv);
2586 program->isSpirV = true;
2587
2588 const auto buildResult{program->rebuildProgramFromIr()};
2589 ASSERT_EQ(CL_SUCCESS, buildResult);
2590
2591 const std::string buildLog{program->getBuildLog(pClDevice->getRootDeviceIndex())};
2592 const auto containsWarning{buildLog.find(CompilerWarnings::recompiledFromIr.data()) != std::string::npos};
2593
2594 EXPECT_TRUE(containsWarning);
2595 }
2596
TEST_F(ProgramTests,givenProgramWithSpirvWhenRebuildIsCalledButSuppressFlagIsEnabledThenRebuildWarningIsNotIssued)2597 TEST_F(ProgramTests, givenProgramWithSpirvWhenRebuildIsCalledButSuppressFlagIsEnabledThenRebuildWarningIsNotIssued) {
2598 const auto program{clUniquePtr(new MockProgram(toClDeviceVector(*pClDevice)))};
2599 uint32_t spirv[16] = {0x03022307, 0x23471113, 0x17192329};
2600 program->irBinary = makeCopy(spirv, sizeof(spirv));
2601 program->irBinarySize = sizeof(spirv);
2602 program->isSpirV = true;
2603
2604 const auto buildOptions{CompilerOptions::noRecompiledFromIr};
2605 program->setBuildOptions(buildOptions.data());
2606
2607 const auto buildResult{program->rebuildProgramFromIr()};
2608 ASSERT_EQ(CL_SUCCESS, buildResult);
2609
2610 const std::string buildLog{program->getBuildLog(pClDevice->getRootDeviceIndex())};
2611 const auto containsWarning{buildLog.find(CompilerWarnings::recompiledFromIr.data()) != std::string::npos};
2612
2613 EXPECT_FALSE(containsWarning);
2614 }
2615
TEST_F(ProgramTests,givenProgramWithSpirvWhenRecompileIsCalledThenRebuildWarningIsIssued)2616 TEST_F(ProgramTests, givenProgramWithSpirvWhenRecompileIsCalledThenRebuildWarningIsIssued) {
2617 const auto program{clUniquePtr(new MockProgram(toClDeviceVector(*pClDevice)))};
2618 uint32_t spirv[16] = {0x03022307, 0x23471113, 0x17192329};
2619 program->irBinary = makeCopy(spirv, sizeof(spirv));
2620 program->irBinarySize = sizeof(spirv);
2621 program->isSpirV = true;
2622
2623 const auto compileResult{program->recompile()};
2624 ASSERT_EQ(CL_SUCCESS, compileResult);
2625
2626 const std::string buildLog{program->getBuildLog(pClDevice->getRootDeviceIndex())};
2627 const auto containsWarning{buildLog.find(CompilerWarnings::recompiledFromIr.data()) != std::string::npos};
2628
2629 EXPECT_TRUE(containsWarning);
2630 }
2631
TEST_F(ProgramTests,givenProgramWithSpirvWhenRecompileIsCalledButSuppressFlagIsEnabledThenRebuildWarningIsNotIssued)2632 TEST_F(ProgramTests, givenProgramWithSpirvWhenRecompileIsCalledButSuppressFlagIsEnabledThenRebuildWarningIsNotIssued) {
2633 const auto program{clUniquePtr(new MockProgram(toClDeviceVector(*pClDevice)))};
2634 uint32_t spirv[16] = {0x03022307, 0x23471113, 0x17192329};
2635 program->irBinary = makeCopy(spirv, sizeof(spirv));
2636 program->irBinarySize = sizeof(spirv);
2637 program->isSpirV = true;
2638
2639 const auto buildOptions{CompilerOptions::noRecompiledFromIr};
2640 program->setBuildOptions(buildOptions.data());
2641
2642 const auto compileResult{program->recompile()};
2643 ASSERT_EQ(CL_SUCCESS, compileResult);
2644
2645 const std::string buildLog{program->getBuildLog(pClDevice->getRootDeviceIndex())};
2646 const auto containsWarning{buildLog.find(CompilerWarnings::recompiledFromIr.data()) != std::string::npos};
2647
2648 EXPECT_FALSE(containsWarning);
2649 }
2650
TEST_F(ProgramTests,whenRebuildingProgramThenStoreDeviceBinaryProperly)2651 TEST_F(ProgramTests, whenRebuildingProgramThenStoreDeviceBinaryProperly) {
2652 auto compilerInterface = new MockCompilerInterface();
2653 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(compilerInterface);
2654 auto compilerMain = new MockCIFMain();
2655 compilerInterface->setIgcMain(compilerMain);
2656 compilerMain->setDefaultCreatorFunc<NEO::MockIgcOclDeviceCtx>(NEO::MockIgcOclDeviceCtx::Create);
2657
2658 MockCompilerDebugVars debugVars = {};
2659 char binaryToReturn[] = "abcdfghijklmnop";
2660 debugVars.binaryToReturn = binaryToReturn;
2661 debugVars.binaryToReturnSize = sizeof(binaryToReturn);
2662 gEnvironment->igcPushDebugVars(debugVars);
2663 std::unique_ptr<void, void (*)(void *)> igcDebugVarsAutoPop{&gEnvironment, [](void *) { gEnvironment->igcPopDebugVars(); }};
2664
2665 auto program = clUniquePtr(new MockProgram(toClDeviceVector(*pClDevice)));
2666 uint32_t ir[16] = {0x03022307, 0x23471113, 0x17192329};
2667 program->irBinary = makeCopy(ir, sizeof(ir));
2668 program->irBinarySize = sizeof(ir);
2669 EXPECT_EQ(nullptr, program->buildInfos[rootDeviceIndex].unpackedDeviceBinary);
2670 EXPECT_EQ(0U, program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize);
2671 program->rebuildProgramFromIr();
2672 ASSERT_NE(nullptr, program->buildInfos[rootDeviceIndex].unpackedDeviceBinary);
2673 ASSERT_EQ(sizeof(binaryToReturn), program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize);
2674 EXPECT_EQ(0, memcmp(binaryToReturn, program->buildInfos[rootDeviceIndex].unpackedDeviceBinary.get(), program->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize));
2675 }
2676
TEST_F(ProgramTests,givenProgramWhenInternalOptionsArePassedThenTheyAreAddedToProgramInternalOptions)2677 TEST_F(ProgramTests, givenProgramWhenInternalOptionsArePassedThenTheyAreAddedToProgramInternalOptions) {
2678 MockProgram program(toClDeviceVector(*pClDevice));
2679 std::string buildOptions = NEO::CompilerOptions::gtpinRera.str();
2680 std::string internalOptions;
2681 program.extractInternalOptions(buildOptions, internalOptions);
2682 EXPECT_STREQ(internalOptions.c_str(), NEO::CompilerOptions::gtpinRera.data());
2683 }
2684
TEST_F(ProgramTests,givenProgramWhenUnknownInternalOptionsArePassedThenTheyAreNotAddedToProgramInternalOptions)2685 TEST_F(ProgramTests, givenProgramWhenUnknownInternalOptionsArePassedThenTheyAreNotAddedToProgramInternalOptions) {
2686 MockProgram program(toClDeviceVector(*pClDevice));
2687 const char *internalOption = "-unknown-internal-options-123";
2688 std::string buildOptions(internalOption);
2689 std::string internalOptions;
2690 program.extractInternalOptions(buildOptions, internalOptions);
2691 EXPECT_EQ(0u, internalOptions.length());
2692 }
2693
TEST_F(ProgramTests,givenProgramWhenAllInternalOptionsArePassedMixedWithUnknownInputThenTheyAreParsedCorrectly)2694 TEST_F(ProgramTests, givenProgramWhenAllInternalOptionsArePassedMixedWithUnknownInputThenTheyAreParsedCorrectly) {
2695 MockProgram program(toClDeviceVector(*pClDevice));
2696 std::string buildOptions = CompilerOptions::concatenate("###", CompilerOptions::gtpinRera, "###", CompilerOptions::greaterThan4gbBuffersRequired, "###");
2697 std::string expectedOutput = CompilerOptions::concatenate(CompilerOptions::gtpinRera, CompilerOptions::greaterThan4gbBuffersRequired);
2698 std::string internalOptions;
2699 program.extractInternalOptions(buildOptions, internalOptions);
2700 EXPECT_EQ(expectedOutput, internalOptions);
2701 }
2702
TEST_F(ProgramTests,givenProgramWhenInternalOptionsArePassedWithValidValuesThenTheyAreAddedToProgramInternalOptions)2703 TEST_F(ProgramTests, givenProgramWhenInternalOptionsArePassedWithValidValuesThenTheyAreAddedToProgramInternalOptions) {
2704 MockProgram program(toClDeviceVector(*pClDevice));
2705
2706 program.isFlagOptionOverride = false;
2707 program.isOptionValueValidOverride = true;
2708 std::string buildOptions = CompilerOptions::concatenate(CompilerOptions::gtpinRera, "someValue");
2709
2710 std::string internalOptions;
2711 program.extractInternalOptions(buildOptions, internalOptions);
2712 EXPECT_EQ(buildOptions, internalOptions) << internalOptions;
2713 }
2714
TEST_F(ProgramTests,givenProgramWhenInternalOptionsArePassedWithInvalidValuesThenTheyAreNotAddedToProgramInternalOptions)2715 TEST_F(ProgramTests, givenProgramWhenInternalOptionsArePassedWithInvalidValuesThenTheyAreNotAddedToProgramInternalOptions) {
2716 MockProgram program(toClDeviceVector(*pClDevice));
2717
2718 program.isFlagOptionOverride = false;
2719 std::string buildOptions = CompilerOptions::concatenate(CompilerOptions::gtpinRera, "someValue");
2720 std::string expectedOutput = "";
2721
2722 std::string internalOptions;
2723 program.extractInternalOptions(buildOptions, internalOptions);
2724 EXPECT_EQ(expectedOutput, internalOptions);
2725
2726 program.isOptionValueValidOverride = true;
2727 buildOptions = std::string(CompilerOptions::gtpinRera);
2728 internalOptions.erase();
2729 program.extractInternalOptions(buildOptions, internalOptions);
2730 EXPECT_EQ(expectedOutput, internalOptions);
2731 }
2732
TEST_F(ProgramTests,GivenInjectInternalBuildOptionsWhenBuildingProgramThenInternalOptionsWereAppended)2733 TEST_F(ProgramTests, GivenInjectInternalBuildOptionsWhenBuildingProgramThenInternalOptionsWereAppended) {
2734 DebugManagerStateRestore dbgRestorer;
2735 DebugManager.flags.InjectInternalBuildOptions.set("-abc");
2736
2737 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
2738 auto pDevice = pContext->getDevice(0);
2739 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
2740 auto program = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pDevice));
2741 program->sourceCode = "__kernel mock() {}";
2742 program->createdFrom = Program::CreatedFrom::SOURCE;
2743
2744 cl_int retVal = program->build(program->getDevices(), "", false);
2745 EXPECT_EQ(CL_SUCCESS, retVal);
2746
2747 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, "-abc")) << cip->buildInternalOptions;
2748 }
2749
TEST_F(ProgramTests,GivenInjectInternalBuildOptionsWhenBuildingBuiltInProgramThenInternalOptionsAreNotAppended)2750 TEST_F(ProgramTests, GivenInjectInternalBuildOptionsWhenBuildingBuiltInProgramThenInternalOptionsAreNotAppended) {
2751 DebugManagerStateRestore dbgRestorer;
2752 DebugManager.flags.InjectInternalBuildOptions.set("-abc");
2753
2754 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
2755 auto pDevice = pContext->getDevice(0);
2756 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
2757 auto program = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pDevice));
2758 program->sourceCode = "__kernel mock() {}";
2759 program->createdFrom = Program::CreatedFrom::SOURCE;
2760 program->isBuiltIn = true;
2761
2762 cl_int retVal = program->build(program->getDevices(), "", false);
2763 EXPECT_EQ(CL_SUCCESS, retVal);
2764
2765 EXPECT_FALSE(CompilerOptions::contains(cip->buildInternalOptions, "-abc")) << cip->buildInternalOptions;
2766 }
2767
TEST_F(ProgramTests,GivenInjectInternalBuildOptionsWhenCompilingProgramThenInternalOptionsWereAppended)2768 TEST_F(ProgramTests, GivenInjectInternalBuildOptionsWhenCompilingProgramThenInternalOptionsWereAppended) {
2769 DebugManagerStateRestore dbgRestorer;
2770 DebugManager.flags.InjectInternalBuildOptions.set("-abc");
2771
2772 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
2773 auto pDevice = pContext->getDevice(0);
2774 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
2775 auto program = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pDevice));
2776 program->sourceCode = "__kernel mock() {}";
2777 program->createdFrom = Program::CreatedFrom::SOURCE;
2778
2779 cl_int retVal = program->compile(program->getDevices(), nullptr, 0, nullptr, nullptr);
2780 EXPECT_EQ(CL_SUCCESS, retVal);
2781
2782 EXPECT_TRUE(CompilerOptions::contains(cip->buildInternalOptions, "-abc")) << cip->buildInternalOptions;
2783 }
2784
TEST_F(ProgramTests,GivenInjectInternalBuildOptionsWhenCompilingBuiltInProgramThenInternalOptionsAreNotAppended)2785 TEST_F(ProgramTests, GivenInjectInternalBuildOptionsWhenCompilingBuiltInProgramThenInternalOptionsAreNotAppended) {
2786 DebugManagerStateRestore dbgRestorer;
2787 DebugManager.flags.InjectInternalBuildOptions.set("-abc");
2788
2789 auto cip = new MockCompilerInterfaceCaptureBuildOptions();
2790 auto pDevice = pContext->getDevice(0);
2791 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(cip);
2792 auto program = std::make_unique<SucceedingGenBinaryProgram>(toClDeviceVector(*pDevice));
2793 program->sourceCode = "__kernel mock() {}";
2794 program->createdFrom = Program::CreatedFrom::SOURCE;
2795 program->isBuiltIn = true;
2796
2797 cl_int retVal = program->compile(program->getDevices(), nullptr, 0, nullptr, nullptr);
2798 EXPECT_EQ(CL_SUCCESS, retVal);
2799
2800 EXPECT_FALSE(CompilerOptions::contains(cip->buildInternalOptions, "-abc")) << cip->buildInternalOptions;
2801 }
2802
2803 class AdditionalOptionsMockProgram : public MockProgram {
2804 public:
2805 using MockProgram::MockProgram;
applyAdditionalOptions(std::string & internalOptions)2806 void applyAdditionalOptions(std::string &internalOptions) override {
2807 applyAdditionalOptionsCalled++;
2808 MockProgram::applyAdditionalOptions(internalOptions);
2809 }
2810 uint32_t applyAdditionalOptionsCalled = 0;
2811 };
2812
TEST_F(ProgramTests,givenProgramWhenBuiltThenAdditionalOptionsAreApplied)2813 TEST_F(ProgramTests, givenProgramWhenBuiltThenAdditionalOptionsAreApplied) {
2814 AdditionalOptionsMockProgram program(toClDeviceVector(*pClDevice));
2815
2816 program.build(program.getDevices(), nullptr, false);
2817 EXPECT_EQ(1u, program.applyAdditionalOptionsCalled);
2818 }
2819
TEST(CreateProgramFromBinaryTests,givenBinaryProgramBuiltInWhenKernelRebulildIsForcedThenDeviceBinaryIsNotUsed)2820 TEST(CreateProgramFromBinaryTests, givenBinaryProgramBuiltInWhenKernelRebulildIsForcedThenDeviceBinaryIsNotUsed) {
2821 DebugManagerStateRestore dbgRestorer;
2822 DebugManager.flags.RebuildPrecompiledKernels.set(true);
2823 cl_int retVal = CL_INVALID_BINARY;
2824
2825 PatchTokensTestData::ValidEmptyProgram programTokens;
2826
2827 auto clDevice = std::make_unique<MockClDevice>(MockDevice::createWithNewExecutionEnvironment<MockDevice>(nullptr));
2828 std::unique_ptr<MockProgram> pProgram(Program::createBuiltInFromGenBinary<MockProgram>(nullptr, toClDeviceVector(*clDevice), programTokens.storage.data(), programTokens.storage.size(), &retVal));
2829 ASSERT_NE(nullptr, pProgram.get());
2830 EXPECT_EQ(CL_SUCCESS, retVal);
2831
2832 auto rootDeviceIndex = clDevice->getRootDeviceIndex();
2833 retVal = pProgram->createProgramFromBinary(programTokens.storage.data(), programTokens.storage.size(), *clDevice);
2834 EXPECT_EQ(CL_SUCCESS, retVal);
2835 EXPECT_EQ(nullptr, pProgram->buildInfos[rootDeviceIndex].unpackedDeviceBinary.get());
2836 EXPECT_EQ(0U, pProgram->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize);
2837 EXPECT_EQ(nullptr, pProgram->buildInfos[rootDeviceIndex].packedDeviceBinary);
2838 EXPECT_EQ(0U, pProgram->buildInfos[rootDeviceIndex].packedDeviceBinarySize);
2839 }
2840
TEST(CreateProgramFromBinaryTests,givenBinaryProgramBuiltInWhenKernelRebulildIsForcedThenRebuildWarningIsEnabled)2841 TEST(CreateProgramFromBinaryTests, givenBinaryProgramBuiltInWhenKernelRebulildIsForcedThenRebuildWarningIsEnabled) {
2842 DebugManagerStateRestore dbgRestorer{};
2843 DebugManager.flags.RebuildPrecompiledKernels.set(true);
2844
2845 PatchTokensTestData::ValidEmptyProgram programTokens;
2846 cl_int retVal{CL_INVALID_BINARY};
2847
2848 const auto clDevice = std::make_unique<MockClDevice>(MockDevice::createWithNewExecutionEnvironment<MockDevice>(nullptr));
2849 std::unique_ptr<MockProgram> pProgram(Program::createBuiltInFromGenBinary<MockProgram>(nullptr, toClDeviceVector(*clDevice), programTokens.storage.data(), programTokens.storage.size(), &retVal));
2850 ASSERT_NE(nullptr, pProgram.get());
2851 ASSERT_EQ(CL_SUCCESS, retVal);
2852
2853 retVal = pProgram->createProgramFromBinary(programTokens.storage.data(), programTokens.storage.size(), *clDevice);
2854 ASSERT_EQ(CL_SUCCESS, retVal);
2855
2856 ASSERT_TRUE(pProgram->shouldWarnAboutRebuild);
2857 }
2858
TEST(CreateProgramFromBinaryTests,givenBinaryProgramNotBuiltInWhenBuiltInKernelRebulildIsForcedThenDeviceBinaryIsUsed)2859 TEST(CreateProgramFromBinaryTests, givenBinaryProgramNotBuiltInWhenBuiltInKernelRebulildIsForcedThenDeviceBinaryIsUsed) {
2860 DebugManagerStateRestore dbgRestorer;
2861 DebugManager.flags.RebuildPrecompiledKernels.set(true);
2862 cl_int retVal = CL_INVALID_BINARY;
2863
2864 PatchTokensTestData::ValidEmptyProgram programTokens;
2865 const unsigned char *binaries[] = {programTokens.storage.data()};
2866 size_t lengths[] = {programTokens.storage.size()};
2867 auto clDevice = std::make_unique<MockClDevice>(MockDevice::createWithNewExecutionEnvironment<MockDevice>(nullptr));
2868 std::unique_ptr<MockProgram> pProgram(Program::create<MockProgram>(
2869 nullptr,
2870 toClDeviceVector(*clDevice),
2871 lengths,
2872 binaries,
2873 nullptr,
2874 retVal));
2875 ASSERT_NE(nullptr, pProgram.get());
2876 EXPECT_EQ(CL_SUCCESS, retVal);
2877
2878 auto rootDeviceIndex = clDevice->getRootDeviceIndex();
2879 EXPECT_NE(nullptr, pProgram->buildInfos[rootDeviceIndex].unpackedDeviceBinary.get());
2880 EXPECT_LT(0U, pProgram->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize);
2881 EXPECT_NE(nullptr, pProgram->buildInfos[rootDeviceIndex].packedDeviceBinary);
2882 EXPECT_LT(0U, pProgram->buildInfos[rootDeviceIndex].packedDeviceBinarySize);
2883 }
2884
TEST(CreateProgramFromBinaryTests,givenBinaryProgramWhenKernelRebulildIsNotForcedThenDeviceBinaryIsUsed)2885 TEST(CreateProgramFromBinaryTests, givenBinaryProgramWhenKernelRebulildIsNotForcedThenDeviceBinaryIsUsed) {
2886 cl_int retVal = CL_INVALID_BINARY;
2887
2888 PatchTokensTestData::ValidEmptyProgram programTokens;
2889
2890 auto clDevice = std::make_unique<MockClDevice>(MockDevice::createWithNewExecutionEnvironment<MockDevice>(nullptr));
2891 std::unique_ptr<MockProgram> pProgram(Program::createBuiltInFromGenBinary<MockProgram>(nullptr, toClDeviceVector(*clDevice), programTokens.storage.data(), programTokens.storage.size(), &retVal));
2892 ASSERT_NE(nullptr, pProgram.get());
2893 EXPECT_EQ(CL_SUCCESS, retVal);
2894
2895 auto rootDeviceIndex = clDevice->getRootDeviceIndex();
2896 retVal = pProgram->createProgramFromBinary(programTokens.storage.data(), programTokens.storage.size(), *clDevice);
2897 EXPECT_EQ(CL_SUCCESS, retVal);
2898 EXPECT_NE(nullptr, reinterpret_cast<uint8_t *>(pProgram->buildInfos[rootDeviceIndex].unpackedDeviceBinary.get()));
2899 EXPECT_EQ(programTokens.storage.size(), pProgram->buildInfos[rootDeviceIndex].unpackedDeviceBinarySize);
2900 EXPECT_NE(nullptr, reinterpret_cast<uint8_t *>(pProgram->buildInfos[rootDeviceIndex].packedDeviceBinary.get()));
2901 EXPECT_EQ(programTokens.storage.size(), pProgram->buildInfos[rootDeviceIndex].packedDeviceBinarySize);
2902 }
2903
2904 struct SpecializationConstantProgramMock : public MockProgram {
2905 using MockProgram::MockProgram;
updateSpecializationConstantSpecializationConstantProgramMock2906 cl_int updateSpecializationConstant(cl_uint specId, size_t specSize, const void *specValue) override {
2907 return CL_SUCCESS;
2908 }
2909 };
2910
2911 struct SpecializationConstantCompilerInterfaceMock : public CompilerInterface {
2912 TranslationOutput::ErrorCode retVal = TranslationOutput::ErrorCode::Success;
2913 int counter = 0;
2914 const char *spirV = nullptr;
getSpecConstantsInfoSpecializationConstantCompilerInterfaceMock2915 TranslationOutput::ErrorCode getSpecConstantsInfo(const NEO::Device &device, ArrayRef<const char> srcSpirV, SpecConstantInfo &output) override {
2916 counter++;
2917 spirV = srcSpirV.begin();
2918 return retVal;
2919 }
returnErrorSpecializationConstantCompilerInterfaceMock2920 void returnError() {
2921 retVal = TranslationOutput::ErrorCode::CompilationFailure;
2922 }
2923 };
2924
2925 struct SpecializationConstantRootDeviceEnvironemnt : public RootDeviceEnvironment {
SpecializationConstantRootDeviceEnvironemntSpecializationConstantRootDeviceEnvironemnt2926 SpecializationConstantRootDeviceEnvironemnt(ExecutionEnvironment &executionEnvironment) : RootDeviceEnvironment(executionEnvironment) {
2927 compilerInterface.reset(new SpecializationConstantCompilerInterfaceMock());
2928 }
getCompilerInterfaceSpecializationConstantRootDeviceEnvironemnt2929 CompilerInterface *getCompilerInterface() override {
2930 return compilerInterface.get();
2931 }
2932
initAilConfigurationSpecializationConstantRootDeviceEnvironemnt2933 bool initAilConfiguration() override {
2934 return true;
2935 }
2936 };
2937
2938 struct setProgramSpecializationConstantTests : public ::testing::Test {
setProgramSpecializationConstantTestssetProgramSpecializationConstantTests2939 setProgramSpecializationConstantTests() : device(new MockDevice()) {}
SetUpsetProgramSpecializationConstantTests2940 void SetUp() override {
2941 mockCompiler = new SpecializationConstantCompilerInterfaceMock();
2942 auto rootDeviceEnvironment = device.getExecutionEnvironment()->rootDeviceEnvironments[0].get();
2943 rootDeviceEnvironment->compilerInterface.reset(mockCompiler);
2944 mockProgram.reset(new SpecializationConstantProgramMock(toClDeviceVector(device)));
2945 mockProgram->isSpirV = true;
2946
2947 EXPECT_FALSE(mockProgram->areSpecializationConstantsInitialized);
2948 EXPECT_EQ(0, mockCompiler->counter);
2949 }
2950
2951 SpecializationConstantCompilerInterfaceMock *mockCompiler = nullptr;
2952 std::unique_ptr<SpecializationConstantProgramMock> mockProgram;
2953 MockClDevice device;
2954
2955 int specValue = 1;
2956 };
2957
TEST_F(setProgramSpecializationConstantTests,whenSetProgramSpecializationConstantThenBinarySourceIsUsed)2958 TEST_F(setProgramSpecializationConstantTests, whenSetProgramSpecializationConstantThenBinarySourceIsUsed) {
2959 auto retVal = mockProgram->setProgramSpecializationConstant(1, sizeof(int), &specValue);
2960
2961 EXPECT_EQ(1, mockCompiler->counter);
2962 EXPECT_EQ(CL_SUCCESS, retVal);
2963 EXPECT_TRUE(mockProgram->areSpecializationConstantsInitialized);
2964 EXPECT_EQ(mockProgram->irBinary.get(), mockCompiler->spirV);
2965 }
2966
TEST_F(setProgramSpecializationConstantTests,whenSetProgramSpecializationConstantMultipleTimesThenSpecializationConstantsAreInitializedOnce)2967 TEST_F(setProgramSpecializationConstantTests, whenSetProgramSpecializationConstantMultipleTimesThenSpecializationConstantsAreInitializedOnce) {
2968 auto retVal = mockProgram->setProgramSpecializationConstant(1, sizeof(int), &specValue);
2969
2970 EXPECT_EQ(1, mockCompiler->counter);
2971 EXPECT_EQ(CL_SUCCESS, retVal);
2972 EXPECT_TRUE(mockProgram->areSpecializationConstantsInitialized);
2973
2974 retVal = mockProgram->setProgramSpecializationConstant(1, sizeof(int), &specValue);
2975
2976 EXPECT_EQ(1, mockCompiler->counter);
2977 EXPECT_EQ(CL_SUCCESS, retVal);
2978 EXPECT_TRUE(mockProgram->areSpecializationConstantsInitialized);
2979 }
2980
TEST_F(setProgramSpecializationConstantTests,givenInvalidGetSpecConstantsInfoReturnValueWhenSetProgramSpecializationConstantThenErrorIsReturned)2981 TEST_F(setProgramSpecializationConstantTests, givenInvalidGetSpecConstantsInfoReturnValueWhenSetProgramSpecializationConstantThenErrorIsReturned) {
2982 mockCompiler->returnError();
2983
2984 auto retVal = mockProgram->setProgramSpecializationConstant(1, sizeof(int), &specValue);
2985
2986 EXPECT_EQ(1, mockCompiler->counter);
2987 EXPECT_EQ(CL_INVALID_VALUE, retVal);
2988 EXPECT_FALSE(mockProgram->areSpecializationConstantsInitialized);
2989 }
2990
TEST(setProgramSpecializationConstantTest,givenUninitializedCompilerinterfaceWhenSetProgramSpecializationConstantThenErrorIsReturned)2991 TEST(setProgramSpecializationConstantTest, givenUninitializedCompilerinterfaceWhenSetProgramSpecializationConstantThenErrorIsReturned) {
2992 auto executionEnvironment = new MockExecutionEnvironment();
2993 executionEnvironment->rootDeviceEnvironments[0] = std::make_unique<NoCompilerInterfaceRootDeviceEnvironment>(*executionEnvironment);
2994 executionEnvironment->rootDeviceEnvironments[0]->setHwInfo(defaultHwInfo.get());
2995 MockClDevice mockDevice(new MockDevice{executionEnvironment, 0});
2996 SpecializationConstantProgramMock mockProgram(toClDeviceVector(mockDevice));
2997
2998 mockProgram.isSpirV = true;
2999 int specValue = 1;
3000
3001 auto retVal = mockProgram.setProgramSpecializationConstant(1, sizeof(int), &specValue);
3002 EXPECT_EQ(CL_OUT_OF_HOST_MEMORY, retVal);
3003 }
3004
3005 using ProgramBinTest = Test<ProgramSimpleFixture>;
3006
TEST_F(ProgramBinTest,givenPrintProgramBinaryProcessingTimeSetWhenBuildProgramThenProcessingTimeIsPrinted)3007 TEST_F(ProgramBinTest, givenPrintProgramBinaryProcessingTimeSetWhenBuildProgramThenProcessingTimeIsPrinted) {
3008 DebugManagerStateRestore restorer;
3009 DebugManager.flags.PrintProgramBinaryProcessingTime.set(true);
3010 testing::internal::CaptureStdout();
3011
3012 CreateProgramFromBinary(pContext, pContext->getDevices(), "kernel_data_param");
3013
3014 auto retVal = pProgram->build(
3015 pProgram->getDevices(),
3016 nullptr,
3017 false);
3018
3019 auto output = testing::internal::GetCapturedStdout();
3020 EXPECT_FALSE(output.compare(0, 14, "Elapsed time: "));
3021 EXPECT_EQ(CL_SUCCESS, retVal);
3022 }
3023
3024 struct DebugDataGuard {
3025 DebugDataGuard(const DebugDataGuard &) = delete;
3026 DebugDataGuard(DebugDataGuard &&) = delete;
3027
DebugDataGuardDebugDataGuard3028 DebugDataGuard() {
3029 for (size_t n = 0; n < sizeof(mockDebugData); n++) {
3030 mockDebugData[n] = (char)n;
3031 }
3032
3033 auto vars = NEO::getIgcDebugVars();
3034 vars.debugDataToReturn = mockDebugData;
3035 vars.debugDataToReturnSize = sizeof(mockDebugData);
3036 NEO::setIgcDebugVars(vars);
3037 }
3038
~DebugDataGuardDebugDataGuard3039 ~DebugDataGuard() {
3040 auto vars = NEO::getIgcDebugVars();
3041 vars.debugDataToReturn = nullptr;
3042 vars.debugDataToReturnSize = 0;
3043 NEO::setIgcDebugVars(vars);
3044 }
3045
3046 char mockDebugData[32];
3047 };
3048
TEST_F(ProgramBinTest,GivenBuildWithDebugDataThenBuildDataAvailableViaGetInfo)3049 TEST_F(ProgramBinTest, GivenBuildWithDebugDataThenBuildDataAvailableViaGetInfo) {
3050 DebugDataGuard debugDataGuard;
3051
3052 const char *sourceCode = "__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n";
3053 pProgram = Program::create<MockProgram>(
3054 pContext,
3055 1,
3056 &sourceCode,
3057 &knownSourceSize,
3058 retVal);
3059 retVal = pProgram->build(pProgram->getDevices(), nullptr, false);
3060 EXPECT_EQ(CL_SUCCESS, retVal);
3061
3062 // Verify
3063 size_t debugDataSize = 0;
3064 retVal = pProgram->getInfo(CL_PROGRAM_DEBUG_INFO_SIZES_INTEL, sizeof(debugDataSize), &debugDataSize, nullptr);
3065 EXPECT_EQ(CL_SUCCESS, retVal);
3066
3067 std::unique_ptr<char[]> debugData{new char[debugDataSize]};
3068 for (size_t n = 0; n < sizeof(debugData); n++) {
3069 debugData[n] = 0;
3070 }
3071 char *pDebugData = &debugData[0];
3072 size_t retData = 0;
3073 bool isOK = true;
3074 retVal = pProgram->getInfo(CL_PROGRAM_DEBUG_INFO_INTEL, 1, &pDebugData, &retData);
3075 EXPECT_EQ(CL_INVALID_VALUE, retVal);
3076 retVal = pProgram->getInfo(CL_PROGRAM_DEBUG_INFO_INTEL, debugDataSize, &pDebugData, &retData);
3077 EXPECT_EQ(CL_SUCCESS, retVal);
3078 cl_uint numDevices;
3079 retVal = clGetProgramInfo(pProgram, CL_PROGRAM_NUM_DEVICES, sizeof(numDevices), &numDevices, nullptr);
3080 EXPECT_EQ(CL_SUCCESS, retVal);
3081 EXPECT_EQ(numDevices * sizeof(debugData), retData);
3082 // Check integrity of returned debug data
3083 for (size_t n = 0; n < debugDataSize; n++) {
3084 if (debugData[n] != (char)n) {
3085 isOK = false;
3086 break;
3087 }
3088 }
3089 EXPECT_TRUE(isOK);
3090 for (size_t n = debugDataSize; n < sizeof(debugData); n++) {
3091 if (debugData[n] != (char)0) {
3092 isOK = false;
3093 break;
3094 }
3095 }
3096 EXPECT_TRUE(isOK);
3097
3098 retData = 0;
3099 retVal = pProgram->getInfo(CL_PROGRAM_DEBUG_INFO_INTEL, debugDataSize, nullptr, &retData);
3100 EXPECT_EQ(CL_SUCCESS, retVal);
3101 EXPECT_EQ(numDevices * sizeof(debugData), retData);
3102 }
3103
TEST_F(ProgramBinTest,GivenDebugDataAvailableWhenLinkingProgramThenDebugDataIsStoredInProgram)3104 TEST_F(ProgramBinTest, GivenDebugDataAvailableWhenLinkingProgramThenDebugDataIsStoredInProgram) {
3105 DebugDataGuard debugDataGuard;
3106
3107 const char *sourceCode = "__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n";
3108 pProgram = Program::create<MockProgram>(
3109 pContext,
3110 1,
3111 &sourceCode,
3112 &knownSourceSize,
3113 retVal);
3114
3115 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
3116 EXPECT_EQ(CL_SUCCESS, retVal);
3117
3118 cl_program programToLink = pProgram;
3119 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &programToLink);
3120 EXPECT_EQ(CL_SUCCESS, retVal);
3121
3122 EXPECT_NE(nullptr, pProgram->getDebugData());
3123 }
3124
3125 using ProgramMultiRootDeviceTests = MultiRootDeviceFixture;
3126
TEST_F(ProgramMultiRootDeviceTests,WhenPrivateSurfaceIsCreatedThenItHasCorrectRootDeviceIndex)3127 TEST_F(ProgramMultiRootDeviceTests, WhenPrivateSurfaceIsCreatedThenItHasCorrectRootDeviceIndex) {
3128 auto program = std::make_unique<MockProgram>(context.get(), false, toClDeviceVector(*device1));
3129 auto infoBlock = std::make_unique<MockKernelInfo>();
3130 infoBlock->setPrivateMemory(1000, false, 8, 0, 0);
3131
3132 program->blockKernelManager->addBlockKernelInfo(infoBlock.release());
3133 program->allocateBlockPrivateSurfaces(*device1);
3134
3135 auto privateSurface = program->getBlockKernelManager()->getPrivateSurface(0);
3136 EXPECT_NE(nullptr, privateSurface);
3137 EXPECT_EQ(expectedRootDeviceIndex, privateSurface->getRootDeviceIndex());
3138 }
3139
TEST_F(ProgramMultiRootDeviceTests,WhenProgramIsCreatedThenBuildInfosVectorIsProperlyResized)3140 TEST_F(ProgramMultiRootDeviceTests, WhenProgramIsCreatedThenBuildInfosVectorIsProperlyResized) {
3141 {
3142 ClDeviceVector deviceVector;
3143 deviceVector.push_back(device1);
3144 deviceVector.push_back(device2);
3145
3146 EXPECT_EQ(1u, deviceVector[0]->getRootDeviceIndex());
3147 auto program = std::make_unique<MockProgram>(context.get(), false, deviceVector);
3148
3149 EXPECT_EQ(3u, program->buildInfos.size());
3150 }
3151 {
3152 ClDeviceVector deviceVector;
3153 deviceVector.push_back(device2);
3154 deviceVector.push_back(device1);
3155
3156 EXPECT_EQ(2u, deviceVector[0]->getRootDeviceIndex());
3157 auto program = std::make_unique<MockProgram>(context.get(), false, deviceVector);
3158
3159 EXPECT_EQ(3u, program->buildInfos.size());
3160 }
3161 }
3162
3163 class MockCompilerInterfaceWithGtpinParam : public CompilerInterface {
3164 public:
link(const NEO::Device & device,const TranslationInput & input,TranslationOutput & output)3165 TranslationOutput::ErrorCode link(
3166 const NEO::Device &device,
3167 const TranslationInput &input,
3168 TranslationOutput &output) override {
3169 gtpinInfoPassed = input.GTPinInput;
3170 return CompilerInterface::link(device, input, output);
3171 }
3172 void *gtpinInfoPassed;
3173 };
3174
TEST_F(ProgramBinTest,GivenSourceKernelWhenLinkingProgramThenGtpinInitInfoIsPassed)3175 TEST_F(ProgramBinTest, GivenSourceKernelWhenLinkingProgramThenGtpinInitInfoIsPassed) {
3176 void *pIgcInitPtr = reinterpret_cast<void *>(0x1234);
3177 gtpinSetIgcInit(pIgcInitPtr);
3178 const char *sourceCode = "__kernel void\nCB(\n__global unsigned int* src, __global unsigned int* dst)\n{\nint id = (int)get_global_id(0);\ndst[id] = src[id];\n}\n";
3179 pProgram = Program::create<MockProgram>(
3180 pContext,
3181 1,
3182 &sourceCode,
3183 &knownSourceSize,
3184 retVal);
3185 std::unique_ptr<MockCompilerInterfaceWithGtpinParam> mockCompilerInterface(new MockCompilerInterfaceWithGtpinParam);
3186
3187 retVal = pProgram->compile(pProgram->getDevices(), nullptr, 0, nullptr, nullptr);
3188 EXPECT_EQ(CL_SUCCESS, retVal);
3189 pDevice->getExecutionEnvironment()->rootDeviceEnvironments[pDevice->getRootDeviceIndex()]->compilerInterface.reset(mockCompilerInterface.get());
3190
3191 cl_program programToLink = pProgram;
3192 retVal = pProgram->link(pProgram->getDevices(), nullptr, 1, &programToLink);
3193
3194 EXPECT_EQ(pIgcInitPtr, mockCompilerInterface->gtpinInfoPassed);
3195 mockCompilerInterface.release();
3196 }
3197
TEST(ProgramReplaceDeviceBinary,GivenBinaryZebinThenUseAsBothPackedAndUnpackedBinaryContainer)3198 TEST(ProgramReplaceDeviceBinary, GivenBinaryZebinThenUseAsBothPackedAndUnpackedBinaryContainer) {
3199 ZebinTestData::ValidEmptyProgram zebin;
3200 std::unique_ptr<char[]> src = makeCopy(zebin.storage.data(), zebin.storage.size());
3201 MockContext context;
3202 auto device = context.getDevice(0);
3203 auto rootDeviceIndex = device->getRootDeviceIndex();
3204 MockProgram program{&context, false, toClDeviceVector(*device)};
3205 program.replaceDeviceBinary(std::move(src), zebin.storage.size(), rootDeviceIndex);
3206 ASSERT_EQ(zebin.storage.size(), program.buildInfos[rootDeviceIndex].packedDeviceBinarySize);
3207 ASSERT_EQ(zebin.storage.size(), program.buildInfos[rootDeviceIndex].unpackedDeviceBinarySize);
3208 ASSERT_NE(nullptr, program.buildInfos[rootDeviceIndex].packedDeviceBinary);
3209 ASSERT_NE(nullptr, program.buildInfos[rootDeviceIndex].unpackedDeviceBinary);
3210 EXPECT_EQ(0, memcmp(program.buildInfos[rootDeviceIndex].packedDeviceBinary.get(), zebin.storage.data(), program.buildInfos[rootDeviceIndex].packedDeviceBinarySize));
3211 EXPECT_EQ(0, memcmp(program.buildInfos[rootDeviceIndex].unpackedDeviceBinary.get(), zebin.storage.data(), program.buildInfos[rootDeviceIndex].unpackedDeviceBinarySize));
3212 }
3213
TEST(ProgramCallbackTest,whenFunctionIsNullptrThenUserDataNeedsToBeNullptr)3214 TEST(ProgramCallbackTest, whenFunctionIsNullptrThenUserDataNeedsToBeNullptr) {
3215 void *userData = nullptr;
3216 EXPECT_TRUE(Program::isValidCallback(nullptr, nullptr));
3217 EXPECT_FALSE(Program::isValidCallback(nullptr, &userData));
3218 }
3219
callbackFuncProgram(cl_program program,void * userData)3220 void CL_CALLBACK callbackFuncProgram(
3221 cl_program program,
3222 void *userData) {
3223 *reinterpret_cast<bool *>(userData) = true;
3224 }
TEST(ProgramCallbackTest,whenFunctionIsNotNullptrThenUserDataDoesntMatter)3225 TEST(ProgramCallbackTest, whenFunctionIsNotNullptrThenUserDataDoesntMatter) {
3226 void *userData = nullptr;
3227 EXPECT_TRUE(Program::isValidCallback(callbackFuncProgram, nullptr));
3228 EXPECT_TRUE(Program::isValidCallback(callbackFuncProgram, &userData));
3229 }
3230
TEST(ProgramCallbackTest,whenInvokeCallbackIsCalledThenFunctionIsProperlyInvoked)3231 TEST(ProgramCallbackTest, whenInvokeCallbackIsCalledThenFunctionIsProperlyInvoked) {
3232 bool functionCalled = false;
3233 MockContext context;
3234 MockProgram program{&context, false, context.getDevices()};
3235 program.invokeCallback(callbackFuncProgram, &functionCalled);
3236
3237 EXPECT_TRUE(functionCalled);
3238
3239 program.invokeCallback(nullptr, nullptr);
3240 }
3241
TEST(BuildProgramTest,givenMultiDeviceProgramWhenBuildingThenStoreAndProcessBinaryOnlyOncePerRootDevice)3242 TEST(BuildProgramTest, givenMultiDeviceProgramWhenBuildingThenStoreAndProcessBinaryOnlyOncePerRootDevice) {
3243 MockProgram *pProgram = nullptr;
3244 std::unique_ptr<char[]> pSource = nullptr;
3245 size_t sourceSize = 0;
3246 std::string testFile;
3247
3248 KernelBinaryHelper kbHelper("CopyBuffer_simd16");
3249
3250 testFile.append(clFiles);
3251 testFile.append("CopyBuffer_simd16.cl");
3252
3253 pSource = loadDataFromFile(
3254 testFile.c_str(),
3255 sourceSize);
3256
3257 ASSERT_NE(0u, sourceSize);
3258 ASSERT_NE(nullptr, pSource);
3259
3260 const char *sources[1] = {pSource.get()};
3261
3262 MockUnrestrictiveContextMultiGPU context;
3263 cl_int retVal = CL_INVALID_PROGRAM;
3264
3265 pProgram = Program::create<MockProgram>(
3266 &context,
3267 1,
3268 sources,
3269 &sourceSize,
3270 retVal);
3271
3272 EXPECT_NE(nullptr, pProgram);
3273 ASSERT_EQ(CL_SUCCESS, retVal);
3274
3275 cl_build_status buildStatus;
3276 for (const auto &device : context.getDevices()) {
3277 retVal = clGetProgramBuildInfo(pProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL);
3278 EXPECT_EQ(CL_SUCCESS, retVal);
3279 EXPECT_EQ(CL_BUILD_NONE, buildStatus);
3280 }
3281
3282 retVal = clBuildProgram(
3283 pProgram,
3284 0,
3285 nullptr,
3286 nullptr,
3287 nullptr,
3288 nullptr);
3289
3290 for (auto &rootDeviceIndex : context.getRootDeviceIndices()) {
3291 EXPECT_EQ(1, pProgram->replaceDeviceBinaryCalledPerRootDevice[rootDeviceIndex]);
3292 EXPECT_EQ(1, pProgram->processGenBinaryCalledPerRootDevice[rootDeviceIndex]);
3293 }
3294 ASSERT_EQ(CL_SUCCESS, retVal);
3295
3296 retVal = clReleaseProgram(pProgram);
3297 EXPECT_EQ(CL_SUCCESS, retVal);
3298 }
3299
TEST(BuildProgramTest,givenMultiDeviceProgramWhenBuildingThenStoreKernelInfoPerEachRootDevice)3300 TEST(BuildProgramTest, givenMultiDeviceProgramWhenBuildingThenStoreKernelInfoPerEachRootDevice) {
3301 MockProgram *pProgram = nullptr;
3302 std::unique_ptr<char[]> pSource = nullptr;
3303 size_t sourceSize = 0;
3304 std::string testFile;
3305
3306 KernelBinaryHelper kbHelper("CopyBuffer_simd16");
3307
3308 testFile.append(clFiles);
3309 testFile.append("CopyBuffer_simd16.cl");
3310
3311 pSource = loadDataFromFile(
3312 testFile.c_str(),
3313 sourceSize);
3314
3315 ASSERT_NE(0u, sourceSize);
3316 ASSERT_NE(nullptr, pSource);
3317
3318 const char *sources[1] = {pSource.get()};
3319
3320 MockUnrestrictiveContextMultiGPU context;
3321 cl_int retVal = CL_INVALID_PROGRAM;
3322
3323 pProgram = Program::create<MockProgram>(
3324 &context,
3325 1,
3326 sources,
3327 &sourceSize,
3328 retVal);
3329
3330 EXPECT_NE(nullptr, pProgram);
3331 ASSERT_EQ(CL_SUCCESS, retVal);
3332
3333 cl_build_status buildStatus;
3334 for (const auto &device : context.getDevices()) {
3335 retVal = clGetProgramBuildInfo(pProgram, device, CL_PROGRAM_BUILD_STATUS, sizeof(buildStatus), &buildStatus, NULL);
3336 EXPECT_EQ(CL_SUCCESS, retVal);
3337 EXPECT_EQ(CL_BUILD_NONE, buildStatus);
3338 }
3339
3340 retVal = clBuildProgram(
3341 pProgram,
3342 0,
3343 nullptr,
3344 nullptr,
3345 nullptr,
3346 nullptr);
3347
3348 ASSERT_EQ(CL_SUCCESS, retVal);
3349 for (auto &rootDeviceIndex : context.getRootDeviceIndices()) {
3350 EXPECT_LT(0u, pProgram->getNumKernels());
3351 for (auto i = 0u; i < pProgram->getNumKernels(); i++) {
3352 EXPECT_NE(nullptr, pProgram->getKernelInfo(i, rootDeviceIndex));
3353 }
3354 }
3355
3356 retVal = clReleaseProgram(pProgram);
3357 EXPECT_EQ(CL_SUCCESS, retVal);
3358 }
3359
TEST(ProgramTest,whenProgramIsBuiltAsAnExecutableForAtLeastOneDeviceThenIsBuiltMethodReturnsTrue)3360 TEST(ProgramTest, whenProgramIsBuiltAsAnExecutableForAtLeastOneDeviceThenIsBuiltMethodReturnsTrue) {
3361 MockSpecializedContext context;
3362 MockProgram program(&context, false, context.getDevices());
3363 EXPECT_FALSE(program.isBuilt());
3364 program.deviceBuildInfos[context.getDevice(0)].buildStatus = CL_BUILD_SUCCESS;
3365 program.deviceBuildInfos[context.getDevice(0)].programBinaryType = CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT;
3366 program.deviceBuildInfos[context.getDevice(1)].buildStatus = CL_BUILD_ERROR;
3367 EXPECT_FALSE(program.isBuilt());
3368 program.deviceBuildInfos[context.getDevice(0)].buildStatus = CL_BUILD_SUCCESS;
3369 program.deviceBuildInfos[context.getDevice(0)].programBinaryType = CL_PROGRAM_BINARY_TYPE_EXECUTABLE;
3370 EXPECT_TRUE(program.isBuilt());
3371 }
3372
TEST(ProgramTest,givenUnlockedProgramWhenRetainForKernelIsCalledThenProgramIsLocked)3373 TEST(ProgramTest, givenUnlockedProgramWhenRetainForKernelIsCalledThenProgramIsLocked) {
3374 MockSpecializedContext context;
3375 MockProgram program(&context, false, context.getDevices());
3376 EXPECT_FALSE(program.isLocked());
3377 program.retainForKernel();
3378 EXPECT_TRUE(program.isLocked());
3379 }
3380
TEST(ProgramTest,givenLockedProgramWhenReleasingForKernelIsCalledForEachRetainThenProgramIsUnlocked)3381 TEST(ProgramTest, givenLockedProgramWhenReleasingForKernelIsCalledForEachRetainThenProgramIsUnlocked) {
3382 MockSpecializedContext context;
3383 MockProgram program(&context, false, context.getDevices());
3384 EXPECT_FALSE(program.isLocked());
3385 program.retainForKernel();
3386 EXPECT_TRUE(program.isLocked());
3387 program.retainForKernel();
3388 EXPECT_TRUE(program.isLocked());
3389 program.releaseForKernel();
3390 EXPECT_TRUE(program.isLocked());
3391 program.releaseForKernel();
3392 EXPECT_FALSE(program.isLocked());
3393 }
3394