1 //
2 // Copyright (c) 2020 The Khronos Group Inc.
3 //
4 // Licensed under the Apache License, Version 2.0 (the "License");
5 // you may not use this file except in compliance with the License.
6 // You may obtain a copy of the License at
7 //
8 // http://www.apache.org/licenses/LICENSE-2.0
9 //
10 // Unless required by applicable law or agreed to in writing, software
11 // distributed under the License is distributed on an "AS IS" BASIS,
12 // WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
13 // See the License for the specific language governing permissions and
14 // limitations under the License.
15 //
16 #include "testBase.h"
17 #include "test_unload_platform_compiler_resources.hpp"
18
19 #include <cassert>
20 #include <chrono>
21 #include <functional>
22 #include <future>
23 #include <initializer_list>
24 #include <stdexcept>
25 #include <string>
26 #include <thread>
27 #include <vector>
28
29 namespace {
30
31 class unload_test_failure : public std::runtime_error {
32 public:
33 using std::runtime_error::runtime_error;
34
unload_test_failure(const std::string & function,cl_int error)35 explicit unload_test_failure(const std::string &function, cl_int error)
36 : std::runtime_error(function + " == " + std::to_string(error))
37 {}
38 };
39
40 class build_base {
41 public:
build_base(cl_context context,cl_device_id device)42 build_base(cl_context context, cl_device_id device)
43 : m_context{ context }, m_device{ device }
44 {}
~build_base()45 virtual ~build_base() { reset(); }
46 build_base(const build_base &) = delete;
47 build_base &operator=(const build_base &) = delete;
48
49 virtual void create() = 0;
50
compile()51 virtual void compile()
52 {
53 assert(nullptr != m_program);
54
55 const cl_int err = clCompileProgram(m_program, 1, &m_device, nullptr, 0,
56 nullptr, nullptr, nullptr, nullptr);
57 if (CL_SUCCESS != err)
58 throw unload_test_failure("clCompileProgram()", err);
59 }
60
link()61 virtual void link()
62 {
63 assert(nullptr != m_program);
64
65 cl_int err = CL_INVALID_PLATFORM;
66 m_executable = clLinkProgram(m_context, 1, &m_device, nullptr, 1,
67 &m_program, nullptr, nullptr, &err);
68 if (CL_SUCCESS != err)
69 throw unload_test_failure("clLinkProgram()", err);
70 if (nullptr == m_executable)
71 throw unload_test_failure("clLinkProgram returned nullptr");
72 }
73
verify()74 virtual void verify()
75 {
76 assert(nullptr != m_executable);
77
78 cl_int err = CL_INVALID_VALUE;
79
80 const clKernelWrapper kernel =
81 clCreateKernel(m_executable, "write_kernel", &err);
82 if (CL_SUCCESS != err)
83 throw unload_test_failure("clCreateKernel()", err);
84
85 const clCommandQueueWrapper queue =
86 clCreateCommandQueue(m_context, m_device, 0, &err);
87 if (CL_SUCCESS != err)
88 throw unload_test_failure("clCreateCommandQueue()", err);
89
90 const clMemWrapper buffer = clCreateBuffer(
91 m_context, CL_MEM_READ_WRITE, sizeof(cl_uint), nullptr, &err);
92 if (CL_SUCCESS != err)
93 throw unload_test_failure("clCreateBuffer()", err);
94
95 cl_uint value = 0;
96
97 err = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
98 if (CL_SUCCESS != err)
99 throw unload_test_failure("clSetKernelArg()", err);
100
101 static const size_t work_size = 1;
102 err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &work_size,
103 nullptr, 0, nullptr, nullptr);
104 if (CL_SUCCESS != err)
105 throw unload_test_failure("clEnqueueNDRangeKernel()", err);
106
107 err = clEnqueueReadBuffer(queue, buffer, CL_BLOCKING, 0,
108 sizeof(cl_uint), &value, 0, nullptr, nullptr);
109 if (CL_SUCCESS != err)
110 throw unload_test_failure("clEnqueueReadBuffer()", err);
111
112 err = clFinish(queue);
113 if (CL_SUCCESS != err) throw unload_test_failure("clFinish()", err);
114
115 if (42 != value)
116 {
117 throw unload_test_failure("Kernel wrote " + std::to_string(value)
118 + ", expected 42");
119 }
120 }
121
reset()122 void reset()
123 {
124 if (m_program)
125 {
126 clReleaseProgram(m_program);
127 m_program = nullptr;
128 }
129 if (m_executable)
130 {
131 clReleaseProgram(m_executable);
132 m_executable = nullptr;
133 }
134 }
135
build()136 void build()
137 {
138 compile();
139 link();
140 }
141
142 protected:
143 const cl_context m_context;
144 const cl_device_id m_device;
145 cl_program m_program{};
146 cl_program m_executable{};
147 };
148
149 /**
150 * @brief initializer_list type for constructing loops over build tests.
151 */
152 using build_list = std::initializer_list<std::reference_wrapper<build_base>>;
153
154 class build_with_source : public build_base {
155 public:
156 using build_base::build_base;
157
create()158 void create() final
159 {
160 assert(nullptr == m_program);
161
162 static const char *sources[] = { write_kernel_source };
163
164 cl_int err = CL_INVALID_PLATFORM;
165 m_program =
166 clCreateProgramWithSource(m_context, 1, sources, nullptr, &err);
167 if (CL_SUCCESS != err)
168 throw unload_test_failure("clCreateProgramWithSource()", err);
169 if (nullptr == m_program)
170 throw unload_test_failure(
171 "clCreateProgramWithSource returned nullptr");
172 }
173 };
174
175 class build_with_binary : public build_base {
176 public:
build_with_binary(const cl_context context,const cl_device_id device,const std::vector<unsigned char> & binary)177 build_with_binary(const cl_context context, const cl_device_id device,
178 const std::vector<unsigned char> &binary)
179 : build_base{ context, device }, m_binary{ binary }
180 {}
181
build_with_binary(const cl_context context,const cl_device_id device)182 build_with_binary(const cl_context context, const cl_device_id device)
183 : build_base{ context, device }
184 {
185 cl_int err = CL_INVALID_VALUE;
186
187 /* Build the program from source */
188 static const char *sources[] = { write_kernel_source };
189 clProgramWrapper program =
190 clCreateProgramWithSource(m_context, 1, sources, nullptr, &err);
191 if (CL_SUCCESS != err)
192 throw unload_test_failure("clCreateProgramWithSource()", err);
193
194 err = clCompileProgram(program, 1, &m_device, nullptr, 0, nullptr,
195 nullptr, nullptr, nullptr);
196 if (CL_SUCCESS != err)
197 throw unload_test_failure("clCompileProgram()", err);
198
199 const clProgramWrapper executable =
200 clLinkProgram(m_context, 1, &m_device, nullptr, 1, &program,
201 nullptr, nullptr, &err);
202 if (CL_SUCCESS != err)
203 throw unload_test_failure("clLinkProgram()", err);
204
205 size_t binary_size;
206 err = clGetProgramInfo(executable, CL_PROGRAM_BINARY_SIZES,
207 sizeof(binary_size), &binary_size, nullptr);
208 if (CL_SUCCESS != err)
209 throw unload_test_failure("clGetProgramInfo()", err);
210
211 m_binary.resize(binary_size);
212
213 /* Grab the program binary */
214 unsigned char *binaries[] = { m_binary.data() };
215 err = clGetProgramInfo(executable, CL_PROGRAM_BINARIES,
216 sizeof(unsigned char *), binaries, nullptr);
217 if (CL_SUCCESS != err)
218 throw unload_test_failure("clGetProgramInfo()", err);
219 }
220
create()221 void create() final
222 {
223 assert(nullptr == m_executable);
224
225 const unsigned char *binaries[] = { m_binary.data() };
226 const size_t binary_sizes[] = { m_binary.size() };
227
228 cl_int err = CL_INVALID_PLATFORM;
229 m_executable = clCreateProgramWithBinary(
230 m_context, 1, &m_device, binary_sizes, binaries, nullptr, &err);
231 if (CL_SUCCESS != err)
232 throw unload_test_failure("clCreateProgramWithBinary()", err);
233 if (nullptr == m_executable)
234 throw unload_test_failure(
235 "clCreateProgramWithBinary returned nullptr");
236 }
237
compile()238 void compile() final
239 {
240 assert(nullptr != m_executable);
241
242 /* Program created from binary, there is nothing to do */
243 }
244
link()245 void link() final
246 {
247 assert(nullptr != m_executable);
248
249 const cl_int err = clBuildProgram(m_executable, 1, &m_device, nullptr,
250 nullptr, nullptr);
251 if (CL_SUCCESS != err)
252 throw unload_test_failure("clBuildProgram()", err);
253 }
254
255 private:
256 std::vector<unsigned char> m_binary;
257 };
258
259 class build_with_il : public build_base {
260 public:
build_with_il(const cl_context context,const cl_platform_id platform,const cl_device_id device)261 build_with_il(const cl_context context, const cl_platform_id platform,
262 const cl_device_id device)
263 : build_base{ context, device }
264 {
265 /* Disable build_with_il if neither core nor extension functionality is
266 * available */
267 m_enabled = false;
268
269 Version version = get_device_cl_version(device);
270 if (version >= Version(2, 1))
271 {
272 std::string sILVersion = get_device_il_version_string(device);
273 if (version < Version(3, 0) || !sILVersion.empty())
274 {
275 m_enabled = true;
276 }
277
278 m_CreateProgramWithIL = clCreateProgramWithIL;
279 }
280 else if (is_extension_available(device, "cl_khr_il_program"))
281 {
282 m_CreateProgramWithIL = (decltype(m_CreateProgramWithIL))
283 clGetExtensionFunctionAddressForPlatform(
284 platform, "clCreateProgramWithILKHR");
285 if (nullptr == m_CreateProgramWithIL)
286 {
287 throw unload_test_failure("cl_khr_il_program supported, but "
288 "function address is nullptr");
289 }
290 m_enabled = true;
291 }
292
293 cl_uint address_bits{};
294 const cl_int err =
295 clGetDeviceInfo(device, CL_DEVICE_ADDRESS_BITS, sizeof(cl_uint),
296 &address_bits, nullptr);
297 if (CL_SUCCESS != err)
298 {
299 throw unload_test_failure("Failure getting device address bits");
300 }
301
302 switch (address_bits)
303 {
304 case 32:
305 m_spirv_binary = write_kernel_32_spv.data();
306 m_spirv_size = write_kernel_32_spv.size();
307 break;
308 case 64:
309 m_spirv_binary = write_kernel_64_spv.data();
310 m_spirv_size = write_kernel_64_spv.size();
311 break;
312 default: throw unload_test_failure("Invalid address bits");
313 }
314 }
315
create()316 void create() final
317 {
318 if (!m_enabled) return;
319
320 assert(nullptr == m_program);
321
322 cl_int err = CL_INVALID_PLATFORM;
323 m_program = m_CreateProgramWithIL(m_context, m_spirv_binary,
324 m_spirv_size, &err);
325 if (CL_SUCCESS != err)
326 throw unload_test_failure("clCreateProgramWithIL()", err);
327 if (nullptr == m_program)
328 throw unload_test_failure("clCreateProgramWithIL returned nullptr");
329 }
330
compile()331 void compile() final
332 {
333 if (!m_enabled) return;
334 build_base::compile();
335 }
336
link()337 void link() final
338 {
339 if (!m_enabled) return;
340 build_base::link();
341 }
342
verify()343 void verify() final
344 {
345 if (!m_enabled) return;
346 build_base::verify();
347 }
348
349 private:
350 void *m_spirv_binary;
351 size_t m_spirv_size;
352 bool m_enabled;
353
354 using CreateProgramWithIL_fn = decltype(&clCreateProgramWithIL);
355 CreateProgramWithIL_fn m_CreateProgramWithIL;
356 };
357 }
358
device_platform(cl_device_id device)359 static cl_platform_id device_platform(cl_device_id device)
360 {
361 cl_platform_id platform;
362 const cl_int err = clGetDeviceInfo(device, CL_DEVICE_PLATFORM,
363 sizeof(platform), &platform, nullptr);
364 if (CL_SUCCESS != err)
365 {
366 log_error("Failure getting platform of tested device\n");
367 return nullptr;
368 }
369
370 return platform;
371 }
372
unload_platform_compiler(const cl_platform_id platform)373 static void unload_platform_compiler(const cl_platform_id platform)
374 {
375 const cl_int err = clUnloadPlatformCompiler(platform);
376 if (CL_SUCCESS != err)
377 throw unload_test_failure("clUnloadPlatformCompiler()", err);
378 }
379
380 /* Test calling the function with a valid platform */
test_unload_valid(cl_device_id device,cl_context,cl_command_queue,int)381 int test_unload_valid(cl_device_id device, cl_context, cl_command_queue, int)
382 {
383 const cl_platform_id platform = device_platform(device);
384 const long int err = clUnloadPlatformCompiler(platform);
385
386 if (CL_SUCCESS != err)
387 {
388 log_error("Test failure: clUnloadPlatformCompiler() == %ld\n", err);
389 return 1;
390 }
391
392 return 0;
393 }
394
395 /* Test calling the function with invalid platform */
test_unload_invalid(cl_device_id,cl_context,cl_command_queue,int)396 int test_unload_invalid(cl_device_id, cl_context, cl_command_queue, int)
397 {
398 const long int err = clUnloadPlatformCompiler(nullptr);
399
400 if (CL_INVALID_PLATFORM != err)
401 {
402 log_error("Test failure: clUnloadPlatformCompiler() == %ld\n", err);
403 return 1;
404 }
405
406 return 0;
407 }
408
409 /* Test calling the function multiple times in a row */
test_unload_repeated(cl_device_id device,cl_context context,cl_command_queue,int)410 int test_unload_repeated(cl_device_id device, cl_context context,
411 cl_command_queue, int)
412 {
413 check_compiler_available(device);
414
415 const cl_platform_id platform = device_platform(device);
416 try
417 {
418 build_with_source source(context, device);
419 build_with_binary binary(context, device);
420 build_with_il il(context, platform, device);
421
422 for (build_base &test : build_list{ source, binary, il })
423 {
424 unload_platform_compiler(platform);
425 unload_platform_compiler(platform);
426
427 test.create();
428 test.build();
429 test.verify();
430 }
431 } catch (const unload_test_failure &e)
432 {
433 log_error("Test failure: %s\n", e.what());
434 return 1;
435 }
436
437 return 0;
438 }
439
440 /* Test calling the function between compilation and linking of programs */
test_unload_compile_unload_link(cl_device_id device,cl_context context,cl_command_queue,int)441 int test_unload_compile_unload_link(cl_device_id device, cl_context context,
442 cl_command_queue, int)
443 {
444 check_compiler_available(device);
445
446 const cl_platform_id platform = device_platform(device);
447 try
448 {
449 build_with_source source(context, device);
450 build_with_binary binary(context, device);
451 build_with_il il(context, platform, device);
452
453 for (build_base &test : build_list{ source, binary, il })
454 {
455 unload_platform_compiler(platform);
456 test.create();
457 test.compile();
458 unload_platform_compiler(platform);
459 test.link();
460 test.verify();
461 }
462 } catch (const unload_test_failure &e)
463 {
464 log_error("Test failure: %s\n", e.what());
465 return 1;
466 }
467
468 return 0;
469 }
470
471 /* Test calling the function between program build and kernel creation */
test_unload_build_unload_create_kernel(cl_device_id device,cl_context context,cl_command_queue,int)472 int test_unload_build_unload_create_kernel(cl_device_id device,
473 cl_context context, cl_command_queue,
474 int)
475 {
476 check_compiler_available(device);
477
478 const cl_platform_id platform = device_platform(device);
479 try
480 {
481 build_with_source source(context, device);
482 build_with_binary binary(context, device);
483 build_with_il il(context, platform, device);
484
485 for (build_base &test : build_list{ source, binary, il })
486 {
487 unload_platform_compiler(platform);
488 test.create();
489 test.build();
490 unload_platform_compiler(platform);
491 test.verify();
492 }
493 } catch (const unload_test_failure &e)
494 {
495 log_error("Test failure: %s\n", e.what());
496 return 1;
497 }
498
499 return 0;
500 }
501
502 /* Test linking together two programs that were built with a call to the unload
503 * function in between */
test_unload_link_different(cl_device_id device,cl_context context,cl_command_queue,int)504 int test_unload_link_different(cl_device_id device, cl_context context,
505 cl_command_queue, int)
506 {
507 check_compiler_available(device);
508
509 const cl_platform_id platform = device_platform(device);
510
511 static const char *sources_1[] = { "unsigned int a() { return 42; }" };
512 static const char *sources_2[] = { R"(
513 unsigned int a();
514 kernel void test(global unsigned int *p)
515 {
516 *p = a();
517 })" };
518
519 cl_int err = CL_INVALID_PLATFORM;
520
521 /* Create and compile program 1 */
522 const clProgramWrapper program_1 =
523 clCreateProgramWithSource(context, 1, sources_1, nullptr, &err);
524 if (CL_SUCCESS != err)
525 {
526 log_error("Test failure: clCreateProgramWithSource() == %ld\n",
527 static_cast<long int>(err));
528 return 1;
529 }
530
531 err = clCompileProgram(program_1, 1, &device, nullptr, 0, nullptr, nullptr,
532 nullptr, nullptr);
533 if (CL_SUCCESS != err)
534 {
535 log_error("Test failure: clCompileProgram() == %ld\n",
536 static_cast<long int>(err));
537 return 1;
538 }
539
540 /* Unload the platform compiler */
541 err = clUnloadPlatformCompiler(platform);
542 if (CL_SUCCESS != err)
543 {
544 log_error("Test failure: clUnloadPlatformCompiler() == %ld\n",
545 static_cast<long int>(err));
546 return 1;
547 }
548
549 /* Create and compile program 2 with the new compiler context */
550 const clProgramWrapper program_2 =
551 clCreateProgramWithSource(context, 1, sources_2, nullptr, &err);
552 if (CL_SUCCESS != err)
553 {
554 log_error("Test failure: clCreateProgramWithSource() == %ld\n",
555 static_cast<long int>(err));
556 return 1;
557 }
558
559 err = clCompileProgram(program_2, 1, &device, nullptr, 0, nullptr, nullptr,
560 nullptr, nullptr);
561 if (CL_SUCCESS != err)
562 {
563 log_error("Test failure: clCompileProgram() == %ld\n",
564 static_cast<long int>(err));
565 return 1;
566 }
567
568 /* Link the two programs into an executable program */
569 const cl_program compiled_programs[] = { program_1, program_2 };
570
571 const clProgramWrapper executable =
572 clLinkProgram(context, 1, &device, nullptr, 2, compiled_programs,
573 nullptr, nullptr, &err);
574 if (CL_SUCCESS != err)
575 {
576 log_error("Test failure: clLinkProgram() == %ld\n",
577 static_cast<long int>(err));
578 return 1;
579 }
580
581 /* Verify execution of a kernel from the linked executable */
582 const clKernelWrapper kernel = clCreateKernel(executable, "test", &err);
583 if (CL_SUCCESS != err)
584 {
585 log_error("Test failure: clCreateKernel() == %ld\n",
586 static_cast<long int>(err));
587 return 1;
588 }
589
590 const clCommandQueueWrapper queue =
591 clCreateCommandQueue(context, device, 0, &err);
592 if (CL_SUCCESS != err)
593 {
594 log_error("Test failure: clCreateCommandQueue() == %ld\n",
595 static_cast<long int>(err));
596 return 1;
597 }
598
599 const clMemWrapper buffer = clCreateBuffer(context, CL_MEM_READ_WRITE,
600 sizeof(cl_uint), nullptr, &err);
601 if (CL_SUCCESS != err)
602 {
603 log_error("Test failure: clCreateBuffer() == %ld\n",
604 static_cast<long int>(err));
605 return 1;
606 }
607
608 cl_uint value = 0;
609
610 err = clSetKernelArg(kernel, 0, sizeof(buffer), &buffer);
611 if (CL_SUCCESS != err)
612 {
613 log_error("Test failure: clSetKernelArg() == %ld\n",
614 static_cast<long int>(err));
615 return 1;
616 }
617
618 static const size_t work_size = 1;
619 err = clEnqueueNDRangeKernel(queue, kernel, 1, nullptr, &work_size, nullptr,
620 0, nullptr, nullptr);
621 if (CL_SUCCESS != err)
622 {
623 log_error("Test failure: clEnqueueNDRangeKernel() == %ld\n",
624 static_cast<long int>(err));
625 return 1;
626 }
627
628 err = clEnqueueReadBuffer(queue, buffer, CL_BLOCKING, 0, sizeof(cl_uint),
629 &value, 0, nullptr, nullptr);
630 if (CL_SUCCESS != err)
631 {
632 log_error("Test failure: clEnqueueReadBuffer() == %ld\n",
633 static_cast<long int>(err));
634 return 1;
635 }
636
637 err = clFinish(queue);
638 if (CL_SUCCESS != err) throw unload_test_failure("clFinish()", err);
639
640 if (42 != value)
641 {
642 log_error("Test failure: Kernel wrote %lu, expected 42)\n",
643 static_cast<long unsigned>(value));
644 return 1;
645 }
646
647 return 0;
648 }
649
650 /* Test calling the function in a thread while others threads are building
651 * programs */
test_unload_build_threaded(cl_device_id device,cl_context context,cl_command_queue,int)652 int test_unload_build_threaded(cl_device_id device, cl_context context,
653 cl_command_queue, int)
654 {
655 using clock = std::chrono::steady_clock;
656
657 check_compiler_available(device);
658
659 const cl_platform_id platform = device_platform(device);
660
661 const auto end = clock::now() + std::chrono::seconds(5);
662
663 const auto unload_thread = [&end, platform] {
664 bool success = true;
665
666 /* Repeatedly unload the compiler */
667 try
668 {
669 while (clock::now() < end)
670 {
671 unload_platform_compiler(platform);
672 }
673 } catch (const unload_test_failure &e)
674 {
675 log_error("Test failure: %s\n", e.what());
676 success = false;
677 }
678
679 return success;
680 };
681
682 const auto build_thread = [&end](build_base *build) {
683 bool success = true;
684
685 try
686 {
687 while (clock::now() < end)
688 {
689 build->create();
690 build->build();
691 build->verify();
692 build->reset();
693 }
694 } catch (unload_test_failure &e)
695 {
696 log_error("Test failure: %s\n", e.what());
697 success = false;
698 }
699
700 return success;
701 };
702
703 build_with_source build_source(context, device);
704 build_with_binary build_binary(context, device);
705 build_with_il build_il(context, platform, device);
706
707 /* Run all threads in parallel and wait for them to finish */
708 std::future<bool> unload_result =
709 std::async(std::launch::async, unload_thread);
710 std::future<bool> build_source_result =
711 std::async(std::launch::async, build_thread, &build_source);
712 std::future<bool> build_binary_result =
713 std::async(std::launch::async, build_thread, &build_binary);
714 std::future<bool> build_il_result =
715 std::async(std::launch::async, build_thread, &build_il);
716
717 bool success = true;
718 if (!unload_result.get())
719 {
720 log_error("unload_thread failed\n");
721 success = false;
722 }
723 if (!build_source_result.get())
724 {
725 log_error("build_with_source failed\n");
726 success = false;
727 }
728 if (!build_binary_result.get())
729 {
730 log_error("build_with_binary failed\n");
731 success = false;
732 }
733 if (!build_il_result.get())
734 {
735 log_error("build_with_il failed\n");
736 success = false;
737 }
738
739 return success ? 0 : 1;
740 }
741
742 /* Test grabbing program build information after calling the unload function */
test_unload_build_info(cl_device_id device,cl_context context,cl_command_queue,int)743 int test_unload_build_info(cl_device_id device, cl_context context,
744 cl_command_queue, int)
745 {
746 check_compiler_available(device);
747
748 const cl_platform_id platform = device_platform(device);
749
750 static const char *sources[] = { write_kernel_source };
751
752 cl_int err = CL_INVALID_PLATFORM;
753 /* Create and build the initial program from source */
754 const clProgramWrapper program =
755 clCreateProgramWithSource(context, 1, sources, nullptr, &err);
756 if (CL_SUCCESS != err)
757 {
758 log_error("Test failure: clCreateProgramWithSource() == %ld\n",
759 static_cast<long int>(err));
760 return 1;
761 }
762
763 static const std::string options("-Dtest");
764
765 err =
766 clBuildProgram(program, 1, &device, options.c_str(), nullptr, nullptr);
767 if (CL_SUCCESS != err)
768 {
769 log_error("Test failure: clCompileProgram() == %ld\n",
770 static_cast<long int>(err));
771 return 1;
772 }
773
774 /* Unload the compiler */
775 err = clUnloadPlatformCompiler(platform);
776 if (CL_SUCCESS != err)
777 {
778 log_error("Test failure: clUnloadPlatformCompiler() == %ld\n",
779 static_cast<long int>(err));
780 return 1;
781 }
782
783 std::vector<cl_program_build_info> infos{ CL_PROGRAM_BUILD_STATUS,
784 CL_PROGRAM_BUILD_OPTIONS,
785 CL_PROGRAM_BUILD_LOG,
786 CL_PROGRAM_BINARY_TYPE };
787
788 if (get_device_cl_version(device) >= Version(2, 0))
789 {
790 infos.push_back(CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE);
791 }
792
793 /* Try grabbing the infos after the compiler unload */
794 for (cl_program_build_info info : infos)
795 {
796 size_t info_size = 0;
797 err = clGetProgramBuildInfo(program, device, info, 0, nullptr,
798 &info_size);
799 if (CL_SUCCESS != err)
800 {
801 log_error("Test failure: clGetProgramBuildInfo() == %ld\n",
802 static_cast<long int>(err));
803 return 1;
804 }
805
806 std::vector<char> info_value(info_size);
807
808 size_t written_size = 0;
809 err = clGetProgramBuildInfo(program, device, info, info_size,
810 &info_value[0], &written_size);
811 if (CL_SUCCESS != err)
812 {
813 log_error("Test failure: clGetProgramBuildInfo() == %ld\n",
814 static_cast<long int>(err));
815 return 1;
816 }
817 else if (written_size != info_size)
818 {
819 log_error("Test failure: Written info value size (%zu) was "
820 "different from "
821 "queried size (%zu).\n",
822 written_size, info_size);
823 return 1;
824 }
825
826 /* Verify the information we know the answer to */
827 switch (info)
828 {
829 case CL_PROGRAM_BUILD_STATUS: {
830 constexpr size_t value_size = sizeof(cl_build_status);
831 if (value_size != info_size)
832 {
833 log_error("Test failure: Expected CL_PROGRAM_BUILD_STATUS "
834 "of size %zu, "
835 "but got %zu\n",
836 value_size, info_size);
837 return 1;
838 }
839 cl_build_status value;
840 memcpy(&value, &info_value[0], value_size);
841 if (CL_BUILD_SUCCESS != value)
842 {
843 log_error(
844 "Test failure: CL_PROGRAM_BUILD_STATUS did not return "
845 "CL_BUILD_SUCCESS (%ld), but %ld\n",
846 static_cast<long int>(CL_BUILD_SUCCESS),
847 static_cast<long int>(value));
848 return 1;
849 }
850 }
851 break;
852
853 case CL_PROGRAM_BUILD_OPTIONS: {
854 const size_t value_size = options.length() + 1;
855 if (value_size != info_size)
856 {
857 log_error("Test failure: Expected CL_PROGRAM_BUILD_OPTIONS "
858 "of size "
859 "%zu, but got %zu\n",
860 value_size, info_size);
861 return 1;
862 }
863 else if (options != &info_value[0])
864 {
865 log_error("Test failure: CL_PROGRAM_BUILD_OPTIONS returned "
866 "\"%s\" "
867 "instead of \"%s\"\n",
868 &info_value[0], options.c_str());
869 return 1;
870 }
871 }
872 break;
873
874 case CL_PROGRAM_BINARY_TYPE: {
875 constexpr size_t value_size = sizeof(cl_program_binary_type);
876 if (value_size != info_size)
877 {
878 log_error("Test failure: Expected CL_PROGRAM_BINARY_TYPE "
879 "of size %zu, "
880 "but got %zu\n",
881 value_size, info_size);
882 return 1;
883 }
884 cl_program_binary_type value;
885 memcpy(&value, &info_value[0], value_size);
886 if (CL_PROGRAM_BINARY_TYPE_EXECUTABLE != value)
887 {
888 log_error(
889 "Test failure: CL_PROGRAM_BINARY_TYPE did not return "
890 "CL_PROGRAM_BINARY_TYPE_EXECUTABLE (%ld), but %ld\n",
891 static_cast<long int>(
892 CL_PROGRAM_BINARY_TYPE_EXECUTABLE),
893 static_cast<long int>(value));
894 return 1;
895 }
896 }
897 break;
898 }
899 }
900
901 return 0;
902 }
903
904 /* Test calling the unload function between program building and fetching the
905 * program binaries */
test_unload_program_binaries(cl_device_id device,cl_context context,cl_command_queue,int)906 int test_unload_program_binaries(cl_device_id device, cl_context context,
907 cl_command_queue, int)
908 {
909 check_compiler_available(device);
910
911 const cl_platform_id platform = device_platform(device);
912
913 static const char *sources[] = { write_kernel_source };
914
915 cl_int err = CL_INVALID_PLATFORM;
916 /* Create and build the initial program from source */
917 const clProgramWrapper program =
918 clCreateProgramWithSource(context, 1, sources, nullptr, &err);
919 if (CL_SUCCESS != err)
920 {
921 log_error("Test failure: clCreateProgramWithSource() == %ld\n",
922 static_cast<long int>(err));
923 return 1;
924 }
925
926 err = clBuildProgram(program, 1, &device, nullptr, nullptr, nullptr);
927 if (CL_SUCCESS != err)
928 {
929 log_error("Test failure: clCompileProgram() == %ld\n",
930 static_cast<long int>(err));
931 return 1;
932 }
933
934 /* Unload the compiler */
935 err = clUnloadPlatformCompiler(platform);
936 if (CL_SUCCESS != err)
937 {
938 log_error("Test failure: clUnloadPlatformCompiler() == %ld\n",
939 static_cast<long int>(err));
940 return 1;
941 }
942
943 /* Grab the built executable binary after the compiler unload */
944 size_t binary_size;
945 err = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES,
946 sizeof(binary_size), &binary_size, nullptr);
947 if (CL_SUCCESS != err)
948 {
949 log_error("Test failure: clGetProgramInfo() == %ld\n",
950 static_cast<long int>(err));
951 return 1;
952 }
953
954 std::vector<unsigned char> binary(binary_size);
955
956 unsigned char *binaries[] = { binary.data() };
957 err = clGetProgramInfo(program, CL_PROGRAM_BINARIES,
958 sizeof(unsigned char *), binaries, nullptr);
959 if (CL_SUCCESS != err)
960 {
961 log_error("Test failure: clGetProgramInfo() == %ld\n",
962 static_cast<long int>(err));
963 return 1;
964 }
965
966 /* Create a new program from the binary and test its execution */
967 try
968 {
969 build_with_binary build_binary(context, device, binary);
970 build_binary.create();
971 build_binary.build();
972 build_binary.verify();
973 } catch (unload_test_failure &e)
974 {
975 log_error("Test failure: %s\n", e.what());
976 return 1;
977 }
978
979 return 0;
980 }
981