1 //
2 // Copyright (c) 2017 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 #if defined(_WIN32)
18 #include <time.h>
19 #elif defined(__linux__) || defined(__APPLE__)
20 #include <sys/time.h>
21 #include <unistd.h>
22 #endif
23 #include "harness/conversions.h"
24
25 #define MAX_LINE_SIZE_IN_PROGRAM 1024
26 #define MAX_LOG_SIZE_IN_PROGRAM 2048
27
28 const char *sample_kernel_start =
29 "__kernel void sample_test(__global float *src, __global int *dst)\n"
30 "{\n"
31 " float temp = 0.0f;\n"
32 " int tid = get_global_id(0);\n";
33
34 const char *sample_kernel_end = "}\n";
35
36 const char *sample_kernel_lines[] = { "dst[tid] = src[tid];\n",
37 "dst[tid] = src[tid] * 3.f;\n",
38 "temp = src[tid] / 4.f;\n",
39 "dst[tid] = dot(temp,src[tid]);\n",
40 "dst[tid] = dst[tid] + temp;\n" };
41
42 /* I compile and link therefore I am. Robert Ioffe */
43 /* The following kernels are used in testing Improved Compilation and Linking
44 * feature */
45
46 const char *simple_kernel = "__kernel void\n"
47 "CopyBuffer(\n"
48 " __global float* src,\n"
49 " __global float* dst )\n"
50 "{\n"
51 " int id = (int)get_global_id(0);\n"
52 " dst[id] = src[id];\n"
53 "}\n";
54
55 const char *simple_kernel_with_defines =
56 "__kernel void\n"
57 "CopyBuffer(\n"
58 " __global float* src,\n"
59 " __global float* dst )\n"
60 "{\n"
61 " int id = (int)get_global_id(0);\n"
62 " float temp = src[id] - 42;\n"
63 " dst[id] = FIRST + temp + SECOND;\n"
64 "}\n";
65
66 const char *simple_kernel_template = "__kernel void\n"
67 "CopyBuffer%d(\n"
68 " __global float* src,\n"
69 " __global float* dst )\n"
70 "{\n"
71 " int id = (int)get_global_id(0);\n"
72 " dst[id] = src[id];\n"
73 "}\n";
74
75 const char *composite_kernel_start = "__kernel void\n"
76 "CompositeKernel(\n"
77 " __global float* src,\n"
78 " __global float* dst )\n"
79 "{\n";
80
81 const char *composite_kernel_end = "}\n";
82
83 const char *composite_kernel_template = " CopyBuffer%d(src, dst);\n";
84
85 const char *composite_kernel_extern_template = "extern __kernel void\n"
86 "CopyBuffer%d(\n"
87 " __global float* src,\n"
88 " __global float* dst );\n";
89
90 const char *another_simple_kernel = "extern __kernel void\n"
91 "CopyBuffer(\n"
92 " __global float* src,\n"
93 " __global float* dst );\n"
94 "__kernel void\n"
95 "AnotherCopyBuffer(\n"
96 " __global float* src,\n"
97 " __global float* dst )\n"
98 "{\n"
99 " CopyBuffer(src, dst);\n"
100 "}\n";
101
102 const char *simple_header = "extern __kernel void\n"
103 "CopyBuffer(\n"
104 " __global float* src,\n"
105 " __global float* dst );\n";
106
107 const char *simple_header_name = "simple_header.h";
108
109 const char *another_simple_kernel_with_header = "#include \"simple_header.h\"\n"
110 "__kernel void\n"
111 "AnotherCopyBuffer(\n"
112 " __global float* src,\n"
113 " __global float* dst )\n"
114 "{\n"
115 " CopyBuffer(src, dst);\n"
116 "}\n";
117
118 const char *header_name_templates[4] = { "simple_header%d.h",
119 "foo/simple_header%d.h",
120 "foo/bar/simple_header%d.h",
121 "foo/bar/baz/simple_header%d.h" };
122
123 const char *include_header_name_templates[4] = {
124 "#include \"simple_header%d.h\"\n", "#include \"foo/simple_header%d.h\"\n",
125 "#include \"foo/bar/simple_header%d.h\"\n",
126 "#include \"foo/bar/baz/simple_header%d.h\"\n"
127 };
128
129 const char *compile_extern_var = "extern constant float foo;\n";
130 const char *compile_extern_struct = "extern constant struct bar bart;\n";
131 const char *compile_extern_function = "extern int baz(int, int);\n";
132
133 const char *compile_static_var = "static constant float foo = 2.78;\n";
134 const char *compile_static_struct = "static constant struct bar {float x, y, "
135 "z, r; int color; } foo = {3.14159};\n";
136 const char *compile_static_function =
137 "static int foo(int x, int y) { return x*x + y*y; }\n";
138
139 const char *compile_regular_var = "constant float foo = 4.0f;\n";
140 const char *compile_regular_struct =
141 "constant struct bar {float x, y, z, r; int color; } foo = {0.f, 0.f, 0.f, "
142 "0.f, 0};\n";
143 const char *compile_regular_function =
144 "int foo(int x, int y) { return x*x + y*y; }\n";
145
146 const char *link_static_var_access = // use with compile_static_var
147 "extern constant float foo;\n"
148 "float access_foo() { return foo; }\n";
149
150 const char *link_static_struct_access = // use with compile_static_struct
151 "extern constant struct bar{float x, y, z, r; int color; } foo;\n"
152 "struct bar access_foo() {return foo; }\n";
153
154 const char *link_static_function_access = // use with compile_static_function
155 "extern int foo(int, int);\n"
156 "int access_foo() { int blah = foo(3, 4); return blah + 5; }\n";
157
test_large_single_compile(cl_context context,cl_device_id deviceID,unsigned int numLines)158 int test_large_single_compile(cl_context context, cl_device_id deviceID,
159 unsigned int numLines)
160 {
161 int error;
162 cl_program program;
163 const char **lines;
164 unsigned int numChoices, i;
165 MTdata d;
166
167 /* First, allocate the array for our line pointers */
168 lines = (const char **)malloc(numLines * sizeof(const char *));
169 if (lines == NULL)
170 {
171 log_error(
172 "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n",
173 numLines, __FILE__, __LINE__);
174 return -1;
175 }
176
177 /* First and last lines are easy */
178 lines[0] = sample_kernel_start;
179 lines[numLines - 1] = sample_kernel_end;
180
181 numChoices = sizeof(sample_kernel_lines) / sizeof(sample_kernel_lines[0]);
182
183 /* Fill the rest with random lines to hopefully prevent much optimization */
184 d = init_genrand(gRandomSeed);
185 for (i = 1; i < numLines - 1; i++)
186 {
187 lines[i] = sample_kernel_lines[genrand_int32(d) % numChoices];
188 }
189 free_mtdata(d);
190 d = NULL;
191
192 /* Try to create a program with these lines */
193 error = create_single_kernel_helper_create_program(context, &program,
194 numLines, lines);
195 if (program == NULL || error != CL_SUCCESS)
196 {
197 log_error("ERROR: Unable to create long test program with %d lines! "
198 "(%s in %s:%d)",
199 numLines, IGetErrorString(error), __FILE__, __LINE__);
200 free(lines);
201 if (program != NULL)
202 {
203 error = clReleaseProgram(program);
204 test_error(error, "Unable to release a program object");
205 }
206 return -1;
207 }
208
209 /* Build it */
210 error = clBuildProgram(program, 1, &deviceID, NULL, NULL, NULL);
211 test_error(error, "Unable to build a long program");
212
213 /* All done! */
214 error = clReleaseProgram(program);
215 test_error(error, "Unable to release a program object");
216
217 free(lines);
218
219 return 0;
220 }
221
test_large_compile(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)222 int test_large_compile(cl_device_id deviceID, cl_context context,
223 cl_command_queue queue, int num_elements)
224 {
225 unsigned int toTest[] = {
226 64, 128, 256, 512, 1024, 2048, 4096, 0
227 }; // 8192, 16384, 32768, 0 };
228 unsigned int i;
229
230 log_info("Testing large compiles...this might take awhile...\n");
231
232 for (i = 0; toTest[i] != 0; i++)
233 {
234 log_info(" %d...\n", toTest[i]);
235
236 #if defined(_WIN32)
237 clock_t start = clock();
238 #elif defined(__linux__) || defined(__APPLE__)
239 timeval time1, time2;
240 gettimeofday(&time1, NULL);
241 #endif
242
243 if (test_large_single_compile(context, deviceID, toTest[i]) != 0)
244 {
245 log_error(
246 "ERROR: long program test failed for %d lines! (in %s:%d)\n",
247 toTest[i], __FILE__, __LINE__);
248 return -1;
249 }
250
251 #if defined(_WIN32)
252 clock_t end = clock();
253 log_perf((float)(end - start) / (float)CLOCKS_PER_SEC, false,
254 "clock() time in secs", "%d lines", toTest[i]);
255 #elif defined(__linux__) || defined(__APPLE__)
256 gettimeofday(&time2, NULL);
257 log_perf((float)(float)(time2.tv_sec - time1.tv_sec)
258 + 1.0e-6 * (time2.tv_usec - time1.tv_usec),
259 false, "wall time in secs", "%d lines", toTest[i]);
260 #endif
261 }
262
263 return 0;
264 }
265
266 static int verifyCopyBuffer(cl_context context, cl_command_queue queue,
267 cl_kernel kernel);
268
269 #if defined(__APPLE__) || defined(__linux)
270 #define _strdup strdup
271 #endif
272
test_large_multi_file_library(cl_context context,cl_device_id deviceID,cl_command_queue queue,unsigned int numLines)273 int test_large_multi_file_library(cl_context context, cl_device_id deviceID,
274 cl_command_queue queue, unsigned int numLines)
275 {
276 int error;
277 cl_program program;
278 cl_program *simple_kernels;
279 const char **lines;
280 unsigned int i;
281 char buffer[MAX_LINE_SIZE_IN_PROGRAM];
282
283 simple_kernels = (cl_program *)malloc(numLines * sizeof(cl_program));
284 if (simple_kernels == NULL)
285 {
286 log_error("ERROR: Unable to allocate kernels array with %d kernels! "
287 "(in %s:%d)\n",
288 numLines, __FILE__, __LINE__);
289 return -1;
290 }
291 /* First, allocate the array for our line pointers */
292 lines = (const char **)malloc((2 * numLines + 2) * sizeof(const char *));
293 if (lines == NULL)
294 {
295 free(simple_kernels);
296 log_error(
297 "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n",
298 (2 * numLines + 2), __FILE__, __LINE__);
299 return -1;
300 }
301
302 for (i = 0; i < numLines; i++)
303 {
304 sprintf(buffer, composite_kernel_extern_template, i);
305 lines[i] = _strdup(buffer);
306 }
307 /* First and last lines are easy */
308 lines[numLines] = composite_kernel_start;
309 lines[2 * numLines + 1] = composite_kernel_end;
310
311 /* Fill the rest with templated kernels */
312 for (i = numLines + 1; i < 2 * numLines + 1; i++)
313 {
314 sprintf(buffer, composite_kernel_template, i - numLines - 1);
315 lines[i] = _strdup(buffer);
316 }
317
318 /* Try to create a program with these lines */
319 error = create_single_kernel_helper_create_program(context, &program,
320 2 * numLines + 2, lines);
321 if (program == NULL || error != CL_SUCCESS)
322 {
323 log_error("ERROR: Unable to create long test program with %d lines! "
324 "(%s) (in %s:%d)\n",
325 numLines, IGetErrorString(error), __FILE__, __LINE__);
326 free(simple_kernels);
327 for (i = 0; i < numLines; i++)
328 {
329 free((void *)lines[i]);
330 free((void *)lines[i + numLines + 1]);
331 }
332 free(lines);
333 if (program != NULL)
334 {
335 error = clReleaseProgram(program);
336 test_error(error, "Unable to release program object");
337 }
338
339 return -1;
340 }
341
342 /* Compile it */
343 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
344 NULL);
345 test_error(error, "Unable to compile a simple program");
346
347 /* Create and compile templated kernels */
348 for (i = 0; i < numLines; i++)
349 {
350 sprintf(buffer, simple_kernel_template, i);
351 const char *kernel_source = _strdup(buffer);
352 simple_kernels[i] =
353 clCreateProgramWithSource(context, 1, &kernel_source, NULL, &error);
354 if (simple_kernels[i] == NULL || error != CL_SUCCESS)
355 {
356 log_error("ERROR: Unable to create long test program with %d "
357 "lines! (%s) (in %s:%d)\n",
358 numLines, IGetErrorString(error), __FILE__, __LINE__);
359 return -1;
360 }
361
362 /* Compile it */
363 error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL,
364 NULL, NULL, NULL);
365 test_error(error, "Unable to compile a simple program");
366
367 free((void *)kernel_source);
368 }
369
370 /* Create library out of compiled templated kernels */
371 cl_program my_newly_minted_library =
372 clLinkProgram(context, 1, &deviceID, "-create-library", numLines,
373 simple_kernels, NULL, NULL, &error);
374 test_error(error, "Unable to create a multi-line library");
375
376 /* Link the program that calls the kernels and the library that contains
377 * them */
378 cl_program programs[2] = { program, my_newly_minted_library };
379 cl_program my_newly_linked_program = clLinkProgram(
380 context, 1, &deviceID, NULL, 2, programs, NULL, NULL, &error);
381 test_error(error, "Unable to link a program with a library");
382
383 // Create the composite kernel
384 cl_kernel kernel =
385 clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
386 test_error(error, "Unable to create a composite kernel");
387
388 // Run the composite kernel and verify the results
389 error = verifyCopyBuffer(context, queue, kernel);
390 if (error != CL_SUCCESS) return error;
391
392 /* All done! */
393 error = clReleaseProgram(program);
394 test_error(error, "Unable to release program object");
395
396 for (i = 0; i < numLines; i++)
397 {
398 free((void *)lines[i]);
399 free((void *)lines[i + numLines + 1]);
400 }
401 free(lines);
402
403 for (i = 0; i < numLines; i++)
404 {
405 error = clReleaseProgram(simple_kernels[i]);
406 test_error(error, "Unable to release program object");
407 }
408 free(simple_kernels);
409
410 error = clReleaseKernel(kernel);
411 test_error(error, "Unable to release kernel object");
412
413 error = clReleaseProgram(my_newly_minted_library);
414 test_error(error, "Unable to release program object");
415
416 error = clReleaseProgram(my_newly_linked_program);
417 test_error(error, "Unable to release program object");
418
419 return 0;
420 }
421
test_multi_file_libraries(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)422 int test_multi_file_libraries(cl_device_id deviceID, cl_context context,
423 cl_command_queue queue, int num_elements)
424 {
425 unsigned int toTest[] = {
426 2, 4, 8, 16, 32, 64, 128, 256, 0
427 }; // 512, 1024, 2048, 4096, 8192, 16384, 32768, 0 };
428 unsigned int i;
429
430 log_info("Testing multi-file libraries ...this might take awhile...\n");
431
432 for (i = 0; toTest[i] != 0; i++)
433 {
434 log_info(" %d...\n", toTest[i]);
435
436 #if defined(_WIN32)
437 clock_t start = clock();
438 #elif defined(__linux__) || defined(__APPLE__)
439 timeval time1, time2;
440 gettimeofday(&time1, NULL);
441 #endif
442
443 if (test_large_multi_file_library(context, deviceID, queue, toTest[i])
444 != 0)
445 {
446 log_error("ERROR: multi-file library program test failed for %d "
447 "lines! (in %s:%d)\n\n",
448 toTest[i], __FILE__, __LINE__);
449 return -1;
450 }
451
452 #if defined(_WIN32)
453 clock_t end = clock();
454 log_perf((float)(end - start) / (float)CLOCKS_PER_SEC, false,
455 "clock() time in secs", "%d lines", toTest[i]);
456 #elif defined(__linux__) || defined(__APPLE__)
457 gettimeofday(&time2, NULL);
458 log_perf((float)(float)(time2.tv_sec - time1.tv_sec)
459 + 1.0e-6 * (time2.tv_usec - time1.tv_usec),
460 false, "wall time in secs", "%d lines", toTest[i]);
461 #endif
462 }
463
464 return 0;
465 }
466
test_large_multiple_embedded_headers(cl_context context,cl_device_id deviceID,cl_command_queue queue,unsigned int numLines)467 int test_large_multiple_embedded_headers(cl_context context,
468 cl_device_id deviceID,
469 cl_command_queue queue,
470 unsigned int numLines)
471 {
472 int error;
473 cl_program program;
474 cl_program *simple_kernels;
475 cl_program *headers;
476 const char **header_names;
477 const char **lines;
478 unsigned int i;
479 char buffer[MAX_LINE_SIZE_IN_PROGRAM];
480
481 simple_kernels = (cl_program *)malloc(numLines * sizeof(cl_program));
482 if (simple_kernels == NULL)
483 {
484 log_error("ERROR: Unable to allocate simple_kernels array with %d "
485 "lines! (in %s:%d)\n",
486 numLines, __FILE__, __LINE__);
487 return -1;
488 }
489 headers = (cl_program *)malloc(numLines * sizeof(cl_program));
490 if (headers == NULL)
491 {
492 log_error("ERROR: Unable to allocate headers array with %d lines! (in "
493 "%s:%d)\n",
494 numLines, __FILE__, __LINE__);
495 return -1;
496 }
497 /* First, allocate the array for our line pointers */
498 header_names = (const char **)malloc(numLines * sizeof(const char *));
499 if (header_names == NULL)
500 {
501 log_error("ERROR: Unable to allocate header_names array with %d lines! "
502 "(in %s:%d)\n",
503 numLines, __FILE__, __LINE__);
504 return -1;
505 }
506 lines = (const char **)malloc((2 * numLines + 2) * sizeof(const char *));
507 if (lines == NULL)
508 {
509 log_error(
510 "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n",
511 (2 * numLines + 2), __FILE__, __LINE__);
512 return -1;
513 }
514
515 for (i = 0; i < numLines; i++)
516 {
517 sprintf(buffer, include_header_name_templates[i % 4], i);
518 lines[i] = _strdup(buffer);
519 sprintf(buffer, header_name_templates[i % 4], i);
520 header_names[i] = _strdup(buffer);
521
522 sprintf(buffer, composite_kernel_extern_template, i);
523 const char *line = buffer;
524 error = create_single_kernel_helper_create_program(context, &headers[i],
525 1, &line);
526 if (headers[i] == NULL || error != CL_SUCCESS)
527 {
528 log_error("ERROR: Unable to create a simple header program! (%s in "
529 "%s:%d)\n",
530 IGetErrorString(error), __FILE__, __LINE__);
531 return -1;
532 }
533 }
534 /* First and last lines are easy */
535 lines[numLines] = composite_kernel_start;
536 lines[2 * numLines + 1] = composite_kernel_end;
537
538 /* Fill the rest with templated kernels */
539 for (i = numLines + 1; i < 2 * numLines + 1; i++)
540 {
541 sprintf(buffer, composite_kernel_template, i - numLines - 1);
542 lines[i] = _strdup(buffer);
543 }
544
545 /* Try to create a program with these lines */
546 error = create_single_kernel_helper_create_program(context, &program,
547 2 * numLines + 2, lines);
548 if (program == NULL || error != CL_SUCCESS)
549 {
550 log_error("ERROR: Unable to create long test program with %d lines! "
551 "(%s) (in %s:%d)\n",
552 numLines, IGetErrorString(error), __FILE__, __LINE__);
553 return -1;
554 }
555
556 /* Compile it */
557 error = clCompileProgram(program, 1, &deviceID, NULL, numLines, headers,
558 header_names, NULL, NULL);
559 test_error(error, "Unable to compile a simple program");
560
561 /* Create and compile templated kernels */
562 for (i = 0; i < numLines; i++)
563 {
564 sprintf(buffer, simple_kernel_template, i);
565 const char *kernel_source = _strdup(buffer);
566 error = create_single_kernel_helper_create_program(
567 context, &simple_kernels[i], 1, &kernel_source);
568 if (simple_kernels[i] == NULL || error != CL_SUCCESS)
569 {
570 log_error("ERROR: Unable to create long test program with %d "
571 "lines! (%s) (in %s:%d)\n",
572 numLines, IGetErrorString(error), __FILE__, __LINE__);
573 return -1;
574 }
575
576 /* Compile it */
577 error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL,
578 NULL, NULL, NULL);
579 test_error(error, "Unable to compile a simple program");
580
581 free((void *)kernel_source);
582 }
583
584 /* Create library out of compiled templated kernels */
585 cl_program my_newly_minted_library =
586 clLinkProgram(context, 1, &deviceID, "-create-library", numLines,
587 simple_kernels, NULL, NULL, &error);
588 test_error(error, "Unable to create a multi-line library");
589
590 /* Link the program that calls the kernels and the library that contains
591 * them */
592 cl_program programs[2] = { program, my_newly_minted_library };
593 cl_program my_newly_linked_program = clLinkProgram(
594 context, 1, &deviceID, NULL, 2, programs, NULL, NULL, &error);
595 test_error(error, "Unable to link a program with a library");
596
597 // Create the composite kernel
598 cl_kernel kernel =
599 clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
600 test_error(error, "Unable to create a composite kernel");
601
602 // Run the composite kernel and verify the results
603 error = verifyCopyBuffer(context, queue, kernel);
604 if (error != CL_SUCCESS) return error;
605
606 /* All done! */
607 error = clReleaseProgram(program);
608 test_error(error, "Unable to release program object");
609
610 for (i = 0; i < numLines; i++)
611 {
612 free((void *)lines[i]);
613 free((void *)header_names[i]);
614 }
615 for (i = numLines + 1; i < 2 * numLines + 1; i++)
616 {
617 free((void *)lines[i]);
618 }
619 free(lines);
620 free(header_names);
621
622 for (i = 0; i < numLines; i++)
623 {
624 error = clReleaseProgram(simple_kernels[i]);
625 test_error(error, "Unable to release program object");
626 error = clReleaseProgram(headers[i]);
627 test_error(error, "Unable to release header program object");
628 }
629 free(simple_kernels);
630 free(headers);
631
632 error = clReleaseKernel(kernel);
633 test_error(error, "Unable to release kernel object");
634
635 error = clReleaseProgram(my_newly_minted_library);
636 test_error(error, "Unable to release program object");
637
638 error = clReleaseProgram(my_newly_linked_program);
639 test_error(error, "Unable to release program object");
640
641 return 0;
642 }
643
test_multiple_embedded_headers(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)644 int test_multiple_embedded_headers(cl_device_id deviceID, cl_context context,
645 cl_command_queue queue, int num_elements)
646 {
647 unsigned int toTest[] = {
648 2, 4, 8, 16, 32, 64, 128, 256, 0
649 }; // 512, 1024, 2048, 4096, 8192, 16384, 32768, 0 };
650 unsigned int i;
651
652 log_info(
653 "Testing multiple embedded headers ...this might take awhile...\n");
654
655 for (i = 0; toTest[i] != 0; i++)
656 {
657 log_info(" %d...\n", toTest[i]);
658
659 #if defined(_WIN32)
660 clock_t start = clock();
661 #elif defined(__linux__) || defined(__APPLE__)
662 timeval time1, time2;
663 gettimeofday(&time1, NULL);
664 #endif
665
666 if (test_large_multiple_embedded_headers(context, deviceID, queue,
667 toTest[i])
668 != 0)
669 {
670 log_error("ERROR: multiple embedded headers program test failed "
671 "for %d lines! (in %s:%d)\n",
672 toTest[i], __FILE__, __LINE__);
673 return -1;
674 }
675
676 #if defined(_WIN32)
677 clock_t end = clock();
678 log_perf((float)(end - start) / (float)CLOCKS_PER_SEC, false,
679 "clock() time in secs", "%d lines", toTest[i]);
680 #elif defined(__linux__) || defined(__APPLE__)
681 gettimeofday(&time2, NULL);
682 log_perf((float)(float)(time2.tv_sec - time1.tv_sec)
683 + 1.0e-6 * (time2.tv_usec - time1.tv_usec),
684 false, "wall time in secs", "%d lines", toTest[i]);
685 #endif
686 }
687
688 return 0;
689 }
690
logbase(double a,double base)691 double logbase(double a, double base) { return log(a) / log(base); }
692
test_large_multiple_libraries(cl_context context,cl_device_id deviceID,cl_command_queue queue,unsigned int numLines)693 int test_large_multiple_libraries(cl_context context, cl_device_id deviceID,
694 cl_command_queue queue, unsigned int numLines)
695 {
696 int error;
697 cl_program *simple_kernels;
698 const char **lines;
699 unsigned int i;
700 char buffer[MAX_LINE_SIZE_IN_PROGRAM];
701 /* I want to create (log2(N)+1)/2 libraries */
702 unsigned int level = (unsigned int)(logbase(numLines, 2.0) + 1.000001) / 2;
703 unsigned int numLibraries = (unsigned int)pow(2.0, level - 1.0);
704 unsigned int numFilesInLib = numLines / numLibraries;
705 cl_program *my_program_and_libraries =
706 (cl_program *)malloc((1 + numLibraries) * sizeof(cl_program));
707 if (my_program_and_libraries == NULL)
708 {
709 log_error("ERROR: Unable to allocate program array with %d programs! "
710 "(in %s:%d)\n",
711 (1 + numLibraries), __FILE__, __LINE__);
712 return -1;
713 }
714
715 log_info("level - %d, numLibraries - %d, numFilesInLib - %d\n", level,
716 numLibraries, numFilesInLib);
717
718 simple_kernels = (cl_program *)malloc(numLines * sizeof(cl_program));
719 if (simple_kernels == NULL)
720 {
721 log_error("ERROR: Unable to allocate kernels array with %d kernels! "
722 "(in %s:%d)\n",
723 numLines, __FILE__, __LINE__);
724 return -1;
725 }
726 /* First, allocate the array for our line pointers */
727 lines = (const char **)malloc((2 * numLines + 2) * sizeof(const char *));
728 if (lines == NULL)
729 {
730 log_error(
731 "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n",
732 (2 * numLines + 2), __FILE__, __LINE__);
733 return -1;
734 }
735
736 for (i = 0; i < numLines; i++)
737 {
738 sprintf(buffer, composite_kernel_extern_template, i);
739 lines[i] = _strdup(buffer);
740 }
741 /* First and last lines are easy */
742 lines[numLines] = composite_kernel_start;
743 lines[2 * numLines + 1] = composite_kernel_end;
744
745 /* Fill the rest with templated kernels */
746 for (i = numLines + 1; i < 2 * numLines + 1; i++)
747 {
748 sprintf(buffer, composite_kernel_template, i - numLines - 1);
749 lines[i] = _strdup(buffer);
750 }
751
752 /* Try to create a program with these lines */
753 error = create_single_kernel_helper_create_program(
754 context, &my_program_and_libraries[0], 2 * numLines + 2, lines);
755 if (my_program_and_libraries[0] == NULL || error != CL_SUCCESS)
756 {
757 log_error("ERROR: Unable to create long test program with %d lines! "
758 "(%s in %s:%d)\n",
759 numLines, IGetErrorString(error), __FILE__, __LINE__);
760 return -1;
761 }
762
763 /* Compile it */
764 error = clCompileProgram(my_program_and_libraries[0], 1, &deviceID, NULL, 0,
765 NULL, NULL, NULL, NULL);
766 test_error(error, "Unable to compile a simple program");
767
768 /* Create and compile templated kernels */
769 for (i = 0; i < numLines; i++)
770 {
771 sprintf(buffer, simple_kernel_template, i);
772 const char *kernel_source = _strdup(buffer);
773 error = create_single_kernel_helper_create_program(
774 context, &simple_kernels[i], 1, &kernel_source);
775 if (simple_kernels[i] == NULL || error != CL_SUCCESS)
776 {
777 log_error("ERROR: Unable to create long test program with %d "
778 "lines! (%s in %s:%d)\n",
779 numLines, IGetErrorString(error), __FILE__, __LINE__);
780 return -1;
781 }
782
783 /* Compile it */
784 error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL,
785 NULL, NULL, NULL);
786 test_error(error, "Unable to compile a simple program");
787
788 free((void *)kernel_source);
789 }
790
791 /* Create library out of compiled templated kernels */
792 for (i = 0; i < numLibraries; i++)
793 {
794 my_program_and_libraries[i + 1] = clLinkProgram(
795 context, 1, &deviceID, "-create-library", numFilesInLib,
796 simple_kernels + i * numFilesInLib, NULL, NULL, &error);
797 test_error(error, "Unable to create a multi-line library");
798 }
799
800 /* Link the program that calls the kernels and the library that contains
801 * them */
802 cl_program my_newly_linked_program =
803 clLinkProgram(context, 1, &deviceID, NULL, numLibraries + 1,
804 my_program_and_libraries, NULL, NULL, &error);
805 test_error(error, "Unable to link a program with a library");
806
807 // Create the composite kernel
808 cl_kernel kernel =
809 clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
810 test_error(error, "Unable to create a composite kernel");
811
812 // Run the composite kernel and verify the results
813 error = verifyCopyBuffer(context, queue, kernel);
814 if (error != CL_SUCCESS) return error;
815
816 /* All done! */
817 for (i = 0; i <= numLibraries; i++)
818 {
819 error = clReleaseProgram(my_program_and_libraries[i]);
820 test_error(error, "Unable to release program object");
821 }
822 free(my_program_and_libraries);
823 for (i = 0; i < numLines; i++)
824 {
825 free((void *)lines[i]);
826 }
827 for (i = numLines + 1; i < 2 * numLines + 1; i++)
828 {
829 free((void *)lines[i]);
830 }
831 free(lines);
832
833 for (i = 0; i < numLines; i++)
834 {
835 error = clReleaseProgram(simple_kernels[i]);
836 test_error(error, "Unable to release program object");
837 }
838 free(simple_kernels);
839
840 error = clReleaseKernel(kernel);
841 test_error(error, "Unable to release kernel object");
842
843 error = clReleaseProgram(my_newly_linked_program);
844 test_error(error, "Unable to release program object");
845
846 return 0;
847 }
848
test_multiple_libraries(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)849 int test_multiple_libraries(cl_device_id deviceID, cl_context context,
850 cl_command_queue queue, int num_elements)
851 {
852 unsigned int toTest[] = {
853 2, 8, 32, 128, 256, 0
854 }; // 512, 2048, 8192, 32768, 0 };
855 unsigned int i;
856
857 log_info("Testing multiple libraries ...this might take awhile...\n");
858
859 for (i = 0; toTest[i] != 0; i++)
860 {
861 log_info(" %d...\n", toTest[i]);
862
863 #if defined(_WIN32)
864 clock_t start = clock();
865 #elif defined(__linux__) || defined(__APPLE__)
866 timeval time1, time2;
867 gettimeofday(&time1, NULL);
868 #endif
869
870 if (test_large_multiple_libraries(context, deviceID, queue, toTest[i])
871 != 0)
872 {
873 log_error("ERROR: multiple library program test failed for %d "
874 "lines! (in %s:%d)\n\n",
875 toTest[i], __FILE__, __LINE__);
876 return -1;
877 }
878
879 #if defined(_WIN32)
880 clock_t end = clock();
881 log_perf((float)(end - start) / (float)CLOCKS_PER_SEC, false,
882 "clock() time in secs", "%d lines", toTest[i]);
883 #elif defined(__linux__) || defined(__APPLE__)
884 gettimeofday(&time2, NULL);
885 log_perf((float)(float)(time2.tv_sec - time1.tv_sec)
886 + 1.0e-6 * (time2.tv_usec - time1.tv_usec),
887 false, "wall time in secs", "%d lines", toTest[i]);
888 #endif
889 }
890
891 return 0;
892 }
893
test_large_multiple_files_multiple_libraries(cl_context context,cl_device_id deviceID,cl_command_queue queue,unsigned int numLines)894 int test_large_multiple_files_multiple_libraries(cl_context context,
895 cl_device_id deviceID,
896 cl_command_queue queue,
897 unsigned int numLines)
898 {
899 int error;
900 cl_program *simple_kernels;
901 const char **lines;
902 unsigned int i;
903 char buffer[MAX_LINE_SIZE_IN_PROGRAM];
904 /* I want to create (log2(N)+1)/4 libraries */
905 unsigned int level = (unsigned int)(logbase(numLines, 2.0) + 1.000001) / 2;
906 unsigned int numLibraries = (unsigned int)pow(2.0, level - 2.0);
907 unsigned int numFilesInLib = numLines / (2 * numLibraries);
908 cl_program *my_programs_and_libraries = (cl_program *)malloc(
909 (1 + numLibraries + numLibraries * numFilesInLib) * sizeof(cl_program));
910 if (my_programs_and_libraries == NULL)
911 {
912 log_error("ERROR: Unable to allocate program array with %d programs! "
913 "(in %s:%d)\n",
914 (1 + numLibraries + numLibraries * numFilesInLib), __FILE__,
915 __LINE__);
916 return -1;
917 }
918 log_info("level - %d, numLibraries - %d, numFilesInLib - %d\n", level,
919 numLibraries, numFilesInLib);
920
921 simple_kernels = (cl_program *)malloc(numLines * sizeof(cl_program));
922 if (simple_kernels == NULL)
923 {
924 log_error("ERROR: Unable to allocate kernels array with %d kernels! "
925 "(in %s:%d)\n",
926 numLines, __FILE__, __LINE__);
927 return -1;
928 }
929 /* First, allocate the array for our line pointers */
930 lines = (const char **)malloc((2 * numLines + 2) * sizeof(const char *));
931 if (lines == NULL)
932 {
933 log_error(
934 "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n",
935 (2 * numLines + 2), __FILE__, __LINE__);
936 return -1;
937 }
938
939 for (i = 0; i < numLines; i++)
940 {
941 sprintf(buffer, composite_kernel_extern_template, i);
942 lines[i] = _strdup(buffer);
943 }
944 /* First and last lines are easy */
945 lines[numLines] = composite_kernel_start;
946 lines[2 * numLines + 1] = composite_kernel_end;
947
948 /* Fill the rest with templated kernels */
949 for (i = numLines + 1; i < 2 * numLines + 1; i++)
950 {
951 sprintf(buffer, composite_kernel_template, i - numLines - 1);
952 lines[i] = _strdup(buffer);
953 }
954
955 /* Try to create a program with these lines */
956 error = create_single_kernel_helper_create_program(
957 context, &my_programs_and_libraries[0], 2 * numLines + 2, lines);
958 if (my_programs_and_libraries[0] == NULL || error != CL_SUCCESS)
959 {
960 log_error("ERROR: Unable to create long test program with %d lines! "
961 "(%s in %s:%d)\n",
962 numLines, IGetErrorString(error), __FILE__, __LINE__);
963 return -1;
964 }
965
966 /* Compile it */
967 error = clCompileProgram(my_programs_and_libraries[0], 1, &deviceID, NULL,
968 0, NULL, NULL, NULL, NULL);
969 test_error(error, "Unable to compile a simple program");
970
971 /* Create and compile templated kernels */
972 for (i = 0; i < numLines; i++)
973 {
974 sprintf(buffer, simple_kernel_template, i);
975 const char *kernel_source = _strdup(buffer);
976 error = create_single_kernel_helper_create_program(
977 context, &simple_kernels[i], 1, &kernel_source);
978 if (simple_kernels[i] == NULL || error != CL_SUCCESS)
979 {
980 log_error("ERROR: Unable to create long test program with %d "
981 "lines! (%s in %s:%d)\n",
982 numLines, IGetErrorString(error), __FILE__, __LINE__);
983 return -1;
984 }
985
986 /* Compile it */
987 error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL,
988 NULL, NULL, NULL);
989 test_error(error, "Unable to compile a simple program");
990
991 free((void *)kernel_source);
992 }
993
994 /* Copy already compiled kernels */
995 for (i = 0; i < numLibraries * numFilesInLib; i++)
996 {
997 my_programs_and_libraries[i + 1] = simple_kernels[i];
998 }
999
1000 /* Create library out of compiled templated kernels */
1001 for (i = 0; i < numLibraries; i++)
1002 {
1003 my_programs_and_libraries[i + 1 + numLibraries * numFilesInLib] =
1004 clLinkProgram(
1005 context, 1, &deviceID, "-create-library", numFilesInLib,
1006 simple_kernels
1007 + (i * numFilesInLib + numLibraries * numFilesInLib),
1008 NULL, NULL, &error);
1009 test_error(error, "Unable to create a multi-line library");
1010 }
1011
1012 /* Link the program that calls the kernels and the library that contains
1013 * them */
1014 cl_program my_newly_linked_program =
1015 clLinkProgram(context, 1, &deviceID, NULL,
1016 numLibraries + 1 + numLibraries * numFilesInLib,
1017 my_programs_and_libraries, NULL, NULL, &error);
1018 test_error(error, "Unable to link a program with a library");
1019
1020 // Create the composite kernel
1021 cl_kernel kernel =
1022 clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
1023 test_error(error, "Unable to create a composite kernel");
1024
1025 // Run the composite kernel and verify the results
1026 error = verifyCopyBuffer(context, queue, kernel);
1027 if (error != CL_SUCCESS) return error;
1028
1029 /* All done! */
1030 for (i = 0; i < numLibraries + 1 + numLibraries * numFilesInLib; i++)
1031 {
1032 error = clReleaseProgram(my_programs_and_libraries[i]);
1033 test_error(error, "Unable to release program object");
1034 }
1035 free(my_programs_and_libraries);
1036
1037 for (i = 0; i < numLines; i++)
1038 {
1039 free((void *)lines[i]);
1040 }
1041 for (i = numLines + 1; i < 2 * numLines + 1; i++)
1042 {
1043 free((void *)lines[i]);
1044 }
1045 free(lines);
1046
1047 for (i = numLibraries * numFilesInLib; i < numLines; i++)
1048 {
1049 error = clReleaseProgram(simple_kernels[i]);
1050 test_error(error, "Unable to release program object");
1051 }
1052 free(simple_kernels);
1053
1054 error = clReleaseKernel(kernel);
1055 test_error(error, "Unable to release kernel object");
1056
1057 error = clReleaseProgram(my_newly_linked_program);
1058 test_error(error, "Unable to release program object");
1059
1060 return 0;
1061 }
1062
test_multiple_files_multiple_libraries(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1063 int test_multiple_files_multiple_libraries(cl_device_id deviceID,
1064 cl_context context,
1065 cl_command_queue queue,
1066 int num_elements)
1067 {
1068 unsigned int toTest[] = { 8, 32, 128, 256,
1069 0 }; // 512, 2048, 8192, 32768, 0 };
1070 unsigned int i;
1071
1072 log_info("Testing multiple files and multiple libraries ...this might take "
1073 "awhile...\n");
1074
1075 for (i = 0; toTest[i] != 0; i++)
1076 {
1077 log_info(" %d...\n", toTest[i]);
1078
1079 #if defined(_WIN32)
1080 clock_t start = clock();
1081 #elif defined(__linux__) || defined(__APPLE__)
1082 timeval time1, time2;
1083 gettimeofday(&time1, NULL);
1084 #endif
1085
1086 if (test_large_multiple_files_multiple_libraries(context, deviceID,
1087 queue, toTest[i])
1088 != 0)
1089 {
1090 log_error("ERROR: multiple files, multiple libraries program test "
1091 "failed for %d lines! (in %s:%d)\n\n",
1092 toTest[i], __FILE__, __LINE__);
1093 return -1;
1094 }
1095
1096 #if defined(_WIN32)
1097 clock_t end = clock();
1098 log_perf((float)(end - start) / (float)CLOCKS_PER_SEC, false,
1099 "clock() time in secs", "%d lines", toTest[i]);
1100 #elif defined(__linux__) || defined(__APPLE__)
1101 gettimeofday(&time2, NULL);
1102 log_perf((float)(float)(time2.tv_sec - time1.tv_sec)
1103 + 1.0e-6 * (time2.tv_usec - time1.tv_usec),
1104 false, "wall time in secs", "%d lines", toTest[i]);
1105 #endif
1106 }
1107
1108 return 0;
1109 }
1110
test_large_multiple_files(cl_context context,cl_device_id deviceID,cl_command_queue queue,unsigned int numLines)1111 int test_large_multiple_files(cl_context context, cl_device_id deviceID,
1112 cl_command_queue queue, unsigned int numLines)
1113 {
1114 int error;
1115 const char **lines;
1116 unsigned int i;
1117 char buffer[MAX_LINE_SIZE_IN_PROGRAM];
1118 cl_program *my_programs =
1119 (cl_program *)malloc((1 + numLines) * sizeof(cl_program));
1120
1121 if (my_programs == NULL)
1122 {
1123 log_error("ERROR: Unable to allocate my_programs array with %d "
1124 "programs! (in %s:%d)\n",
1125 (1 + numLines), __FILE__, __LINE__);
1126 return -1;
1127 }
1128 /* First, allocate the array for our line pointers */
1129 lines = (const char **)malloc((2 * numLines + 2) * sizeof(const char *));
1130 if (lines == NULL)
1131 {
1132 log_error(
1133 "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n",
1134 (2 * numLines + 2), __FILE__, __LINE__);
1135 return -1;
1136 }
1137
1138 for (i = 0; i < numLines; i++)
1139 {
1140 sprintf(buffer, composite_kernel_extern_template, i);
1141 lines[i] = _strdup(buffer);
1142 }
1143 /* First and last lines are easy */
1144 lines[numLines] = composite_kernel_start;
1145 lines[2 * numLines + 1] = composite_kernel_end;
1146
1147 /* Fill the rest with templated kernels */
1148 for (i = numLines + 1; i < 2 * numLines + 1; i++)
1149 {
1150 sprintf(buffer, composite_kernel_template, i - numLines - 1);
1151 lines[i] = _strdup(buffer);
1152 }
1153
1154 /* Try to create a program with these lines */
1155 error = create_single_kernel_helper_create_program(context, &my_programs[0],
1156 2 * numLines + 2, lines);
1157 if (my_programs[0] == NULL || error != CL_SUCCESS)
1158 {
1159 log_error("ERROR: Unable to create long test program with %d lines! "
1160 "(%s in %s:%d)\n",
1161 numLines, IGetErrorString(error), __FILE__, __LINE__);
1162 return -1;
1163 }
1164
1165 /* Compile it */
1166 error = clCompileProgram(my_programs[0], 1, &deviceID, NULL, 0, NULL, NULL,
1167 NULL, NULL);
1168 test_error(error, "Unable to compile a simple program");
1169
1170 /* Create and compile templated kernels */
1171 for (i = 0; i < numLines; i++)
1172 {
1173 sprintf(buffer, simple_kernel_template, i);
1174 const char *kernel_source = _strdup(buffer);
1175 error = create_single_kernel_helper_create_program(
1176 context, &my_programs[i + 1], 1, &kernel_source);
1177 if (my_programs[i + 1] == NULL || error != CL_SUCCESS)
1178 {
1179 log_error("ERROR: Unable to create long test program with %d "
1180 "lines! (%s in %s:%d)\n",
1181 numLines, IGetErrorString(error), __FILE__, __LINE__);
1182 return -1;
1183 }
1184
1185 /* Compile it */
1186 error = clCompileProgram(my_programs[i + 1], 1, &deviceID, NULL, 0,
1187 NULL, NULL, NULL, NULL);
1188 test_error(error, "Unable to compile a simple program");
1189
1190 free((void *)kernel_source);
1191 }
1192
1193 /* Link the program that calls the kernels and the library that contains
1194 * them */
1195 cl_program my_newly_linked_program =
1196 clLinkProgram(context, 1, &deviceID, NULL, 1 + numLines, my_programs,
1197 NULL, NULL, &error);
1198 test_error(error, "Unable to link a program with a library");
1199
1200 // Create the composite kernel
1201 cl_kernel kernel =
1202 clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
1203 test_error(error, "Unable to create a composite kernel");
1204
1205 // Run the composite kernel and verify the results
1206 error = verifyCopyBuffer(context, queue, kernel);
1207 if (error != CL_SUCCESS) return error;
1208
1209 /* All done! */
1210 for (i = 0; i < 1 + numLines; i++)
1211 {
1212 error = clReleaseProgram(my_programs[i]);
1213 test_error(error, "Unable to release program object");
1214 }
1215 free(my_programs);
1216 for (i = 0; i < numLines; i++)
1217 {
1218 free((void *)lines[i]);
1219 }
1220 for (i = numLines + 1; i < 2 * numLines + 1; i++)
1221 {
1222 free((void *)lines[i]);
1223 }
1224 free(lines);
1225
1226 error = clReleaseKernel(kernel);
1227 test_error(error, "Unable to release kernel object");
1228
1229 error = clReleaseProgram(my_newly_linked_program);
1230 test_error(error, "Unable to release program object");
1231
1232 return 0;
1233 }
1234
test_multiple_files(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1235 int test_multiple_files(cl_device_id deviceID, cl_context context,
1236 cl_command_queue queue, int num_elements)
1237 {
1238 unsigned int toTest[] = { 8, 32, 128, 256,
1239 0 }; // 512, 2048, 8192, 32768, 0 };
1240 unsigned int i;
1241
1242 log_info("Testing multiple files compilation and linking into a single "
1243 "executable ...this might take awhile...\n");
1244
1245 for (i = 0; toTest[i] != 0; i++)
1246 {
1247 log_info(" %d...\n", toTest[i]);
1248
1249 #if defined(_WIN32)
1250 clock_t start = clock();
1251 #elif defined(__linux__) || defined(__APPLE__)
1252 timeval time1, time2;
1253 gettimeofday(&time1, NULL);
1254 #endif
1255
1256 if (test_large_multiple_files(context, deviceID, queue, toTest[i]) != 0)
1257 {
1258 log_error("ERROR: multiple files program test failed for %d lines! "
1259 "(in %s:%d)\n\n",
1260 toTest[i], __FILE__, __LINE__);
1261 return -1;
1262 }
1263
1264 #if defined(_WIN32)
1265 clock_t end = clock();
1266 log_perf((float)(end - start) / (float)CLOCKS_PER_SEC, false,
1267 "clock() time in secs", "%d lines", toTest[i]);
1268 #elif defined(__linux__) || defined(__APPLE__)
1269 gettimeofday(&time2, NULL);
1270 log_perf((float)(float)(time2.tv_sec - time1.tv_sec)
1271 + 1.0e-6 * (time2.tv_usec - time1.tv_usec),
1272 false, "wall time in secs", "%d lines", toTest[i]);
1273 #endif
1274 }
1275
1276 return 0;
1277 }
1278
test_simple_compile_only(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1279 int test_simple_compile_only(cl_device_id deviceID, cl_context context,
1280 cl_command_queue queue, int num_elements)
1281 {
1282 int error;
1283 cl_program program;
1284
1285 log_info("Testing a simple compilation only...\n");
1286 error = create_single_kernel_helper_create_program(context, &program, 1,
1287 &simple_kernel);
1288 if (program == NULL || error != CL_SUCCESS)
1289 {
1290 log_error(
1291 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
1292 IGetErrorString(error), __FILE__, __LINE__);
1293 return -1;
1294 }
1295
1296 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1297 NULL);
1298 test_error(error, "Unable to compile a simple program");
1299
1300 /* All done! */
1301 error = clReleaseProgram(program);
1302 test_error(error, "Unable to release program object");
1303
1304 return 0;
1305 }
1306
test_simple_static_compile_only(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1307 int test_simple_static_compile_only(cl_device_id deviceID, cl_context context,
1308 cl_command_queue queue, int num_elements)
1309 {
1310 int error;
1311 cl_program program;
1312
1313 log_info("Testing a simple static compilations only...\n");
1314
1315 error = create_single_kernel_helper_create_program(context, &program, 1,
1316 &compile_static_var);
1317 if (program == NULL || error != CL_SUCCESS)
1318 {
1319 log_error("ERROR: Unable to create a simple static variable test "
1320 "program! (%s in %s:%d)\n",
1321 IGetErrorString(error), __FILE__, __LINE__);
1322 return -1;
1323 }
1324
1325 log_info("Compiling a static variable...\n");
1326 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1327 NULL);
1328 test_error(error, "Unable to compile a simple static variable program");
1329
1330 /* All done! */
1331 error = clReleaseProgram(program);
1332 test_error(error, "Unable to release program object");
1333
1334 error = create_single_kernel_helper_create_program(context, &program, 1,
1335 &compile_static_struct);
1336 if (program == NULL || error != CL_SUCCESS)
1337 {
1338 log_error("ERROR: Unable to create a simple static struct test "
1339 "program! (%s in %s:%d)\n",
1340 IGetErrorString(error), __FILE__, __LINE__);
1341 return -1;
1342 }
1343
1344 log_info("Compiling a static struct...\n");
1345 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1346 NULL);
1347 test_error(error, "Unable to compile a simple static variable program");
1348
1349 /* All done! */
1350 error = clReleaseProgram(program);
1351 test_error(error, "Unable to release program object");
1352
1353 error = create_single_kernel_helper_create_program(
1354 context, &program, 1, &compile_static_function);
1355 if (program == NULL || error != CL_SUCCESS)
1356 {
1357 log_error("ERROR: Unable to create a simple static function test "
1358 "program! (%s in %s:%d)\n",
1359 IGetErrorString(error), __FILE__, __LINE__);
1360 return -1;
1361 }
1362
1363 log_info("Compiling a static function...\n");
1364 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1365 NULL);
1366 test_error(error, "Unable to compile a simple static function program");
1367
1368 /* All done! */
1369 error = clReleaseProgram(program);
1370 test_error(error, "Unable to release program object");
1371
1372 return 0;
1373 }
1374
test_simple_extern_compile_only(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1375 int test_simple_extern_compile_only(cl_device_id deviceID, cl_context context,
1376 cl_command_queue queue, int num_elements)
1377 {
1378 int error;
1379 cl_program program;
1380
1381 log_info("Testing a simple extern compilations only...\n");
1382 error = create_single_kernel_helper_create_program(context, &program, 1,
1383 &simple_header);
1384 if (program == NULL || error != CL_SUCCESS)
1385 {
1386 log_error("ERROR: Unable to create a simple extern kernel test "
1387 "program! (%s in %s:%d)\n",
1388 IGetErrorString(error), __FILE__, __LINE__);
1389 return -1;
1390 }
1391
1392 log_info("Compiling an extern kernel...\n");
1393 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1394 NULL);
1395 test_error(error, "Unable to compile a simple extern kernel program");
1396
1397 /* All done! */
1398 error = clReleaseProgram(program);
1399 test_error(error, "Unable to release program object");
1400
1401 error = create_single_kernel_helper_create_program(context, &program, 1,
1402 &compile_extern_var);
1403 if (program == NULL || error != CL_SUCCESS)
1404 {
1405 log_error("ERROR: Unable to create a simple extern variable test "
1406 "program! (%s in %s:%d)\n",
1407 IGetErrorString(error), __FILE__, __LINE__);
1408 return -1;
1409 }
1410
1411 log_info("Compiling an extern variable...\n");
1412 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1413 NULL);
1414 test_error(error, "Unable to compile a simple extern variable program");
1415
1416 /* All done! */
1417 error = clReleaseProgram(program);
1418 test_error(error, "Unable to release program object");
1419
1420 error = create_single_kernel_helper_create_program(context, &program, 1,
1421 &compile_extern_struct);
1422 if (program == NULL || error != CL_SUCCESS)
1423 {
1424 log_error("ERROR: Unable to create a simple extern struct test "
1425 "program! (%s in %s:%d)\n",
1426 IGetErrorString(error), __FILE__, __LINE__);
1427 return -1;
1428 }
1429
1430 log_info("Compiling an extern struct...\n");
1431 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1432 NULL);
1433 test_error(error, "Unable to compile a simple extern variable program");
1434
1435 /* All done! */
1436 error = clReleaseProgram(program);
1437 test_error(error, "Unable to release program object");
1438
1439 error = create_single_kernel_helper_create_program(
1440 context, &program, 1, &compile_extern_function);
1441 if (program == NULL || error != CL_SUCCESS)
1442 {
1443 log_error("ERROR: Unable to create a simple extern function test "
1444 "program! (%s in %s:%d)\n",
1445 IGetErrorString(error), __FILE__, __LINE__);
1446 return -1;
1447 }
1448
1449 log_info("Compiling an extern function...\n");
1450 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1451 NULL);
1452 test_error(error, "Unable to compile a simple extern function program");
1453
1454 /* All done! */
1455 error = clReleaseProgram(program);
1456 test_error(error, "Unable to release program object");
1457
1458 return 0;
1459 }
1460
1461 struct simple_user_data
1462 {
1463 const char *m_message;
1464 cl_event m_event;
1465 };
1466
1467 const char *once_upon_a_midnight_dreary = "Once upon a midnight dreary!";
1468
simple_compile_callback(cl_program program,void * user_data)1469 static void CL_CALLBACK simple_compile_callback(cl_program program,
1470 void *user_data)
1471 {
1472 simple_user_data *simple_compile_user_data = (simple_user_data *)user_data;
1473 log_info("in the simple_compile_callback: program %p just completed "
1474 "compiling with '%s'\n",
1475 program, simple_compile_user_data->m_message);
1476 if (strcmp(once_upon_a_midnight_dreary, simple_compile_user_data->m_message)
1477 != 0)
1478 {
1479 log_error("ERROR: in the simple_compile_callback: Expected '%s' and "
1480 "got %s (in %s:%d)!\n",
1481 once_upon_a_midnight_dreary,
1482 simple_compile_user_data->m_message, __FILE__, __LINE__);
1483 }
1484
1485 int error;
1486 log_info("in the simple_compile_callback: program %p just completed "
1487 "compiling with '%p'\n",
1488 program, simple_compile_user_data->m_event);
1489
1490 error =
1491 clSetUserEventStatus(simple_compile_user_data->m_event, CL_COMPLETE);
1492 if (error != CL_SUCCESS)
1493 {
1494 log_error("ERROR: in the simple_compile_callback: Unable to set user "
1495 "event status to CL_COMPLETE! (%s in %s:%d)\n",
1496 IGetErrorString(error), __FILE__, __LINE__);
1497 exit(-1);
1498 }
1499 log_info("in the simple_compile_callback: Successfully signaled "
1500 "compile_program_completion_event!\n");
1501 }
1502
test_simple_compile_with_callback(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1503 int test_simple_compile_with_callback(cl_device_id deviceID, cl_context context,
1504 cl_command_queue queue, int num_elements)
1505 {
1506 int error;
1507 cl_program program;
1508 cl_event compile_program_completion_event;
1509
1510 log_info("Testing a simple compilation with callback...\n");
1511 error = create_single_kernel_helper_create_program(context, &program, 1,
1512 &simple_kernel);
1513 if (program == NULL || error != CL_SUCCESS)
1514 {
1515 log_error(
1516 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
1517 IGetErrorString(error), __FILE__, __LINE__);
1518 return -1;
1519 }
1520
1521 compile_program_completion_event = clCreateUserEvent(context, &error);
1522 test_error(error, "Unable to create a user event");
1523
1524 simple_user_data simple_compile_user_data = {
1525 once_upon_a_midnight_dreary, compile_program_completion_event
1526 };
1527
1528 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL,
1529 simple_compile_callback,
1530 (void *)&simple_compile_user_data);
1531 test_error(error, "Unable to compile a simple program with a callback");
1532
1533 error = clWaitForEvents(1, &compile_program_completion_event);
1534 test_error(error,
1535 "clWaitForEvents failed when waiting on "
1536 "compile_program_completion_event");
1537
1538 /* All done! */
1539 error = clReleaseEvent(compile_program_completion_event);
1540 test_error(error, "Unable to release event object");
1541
1542 error = clReleaseProgram(program);
1543 test_error(error, "Unable to release program object");
1544
1545 return 0;
1546 }
1547
test_simple_embedded_header_compile(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1548 int test_simple_embedded_header_compile(cl_device_id deviceID,
1549 cl_context context,
1550 cl_command_queue queue,
1551 int num_elements)
1552 {
1553 int error;
1554 cl_program program, header;
1555
1556 log_info("Testing a simple embedded header compile only...\n");
1557 program = clCreateProgramWithSource(
1558 context, 1, &another_simple_kernel_with_header, NULL, &error);
1559 if (program == NULL || error != CL_SUCCESS)
1560 {
1561 log_error(
1562 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
1563 IGetErrorString(error), __FILE__, __LINE__);
1564 return -1;
1565 }
1566
1567 header =
1568 clCreateProgramWithSource(context, 1, &simple_header, NULL, &error);
1569 if (header == NULL || error != CL_SUCCESS)
1570 {
1571 log_error(
1572 "ERROR: Unable to create a simple header program! (%s in %s:%d)\n",
1573 IGetErrorString(error), __FILE__, __LINE__);
1574 return -1;
1575 }
1576
1577 error = clCompileProgram(program, 1, &deviceID, NULL, 1, &header,
1578 &simple_header_name, NULL, NULL);
1579 test_error(error,
1580 "Unable to compile a simple program with embedded header");
1581
1582 /* All done! */
1583 error = clReleaseProgram(program);
1584 test_error(error, "Unable to release program object");
1585
1586 error = clReleaseProgram(header);
1587 test_error(error, "Unable to release program object");
1588
1589 return 0;
1590 }
1591
test_simple_link_only(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1592 int test_simple_link_only(cl_device_id deviceID, cl_context context,
1593 cl_command_queue queue, int num_elements)
1594 {
1595 int error;
1596 cl_program program;
1597
1598 log_info("Testing a simple linking only...\n");
1599 error = create_single_kernel_helper_create_program(context, &program, 1,
1600 &simple_kernel);
1601 if (program == NULL || error != CL_SUCCESS)
1602 {
1603 log_error(
1604 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
1605 IGetErrorString(error), __FILE__, __LINE__);
1606 return -1;
1607 }
1608
1609 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1610 NULL);
1611 test_error(error, "Unable to compile a simple program");
1612
1613 cl_program my_newly_linked_program = clLinkProgram(
1614 context, 1, &deviceID, NULL, 1, &program, NULL, NULL, &error);
1615 test_error(error, "Unable to link a simple program");
1616
1617 /* All done! */
1618 error = clReleaseProgram(program);
1619 test_error(error, "Unable to release program object");
1620
1621 error = clReleaseProgram(my_newly_linked_program);
1622 test_error(error, "Unable to release program object");
1623
1624 return 0;
1625 }
1626
test_two_file_regular_variable_access(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1627 int test_two_file_regular_variable_access(cl_device_id deviceID,
1628 cl_context context,
1629 cl_command_queue queue,
1630 int num_elements)
1631 {
1632 int error;
1633 cl_program program, second_program, my_newly_linked_program;
1634
1635 const char *sources[2] = {
1636 simple_kernel, compile_regular_var
1637 }; // here we want to avoid linking error due to lack of kernels
1638 log_info("Compiling and linking two program objects, where one tries to "
1639 "access regular variable from another...\n");
1640 error = create_single_kernel_helper_create_program(context, &program, 2,
1641 sources);
1642 if (program == NULL || error != CL_SUCCESS)
1643 {
1644 log_error("ERROR: Unable to create a test program with regular "
1645 "variable! (%s in %s:%d)\n",
1646 IGetErrorString(error), __FILE__, __LINE__);
1647 return -1;
1648 }
1649
1650 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1651 NULL);
1652 test_error(error,
1653 "Unable to compile a simple program with regular function");
1654
1655 error = create_single_kernel_helper_create_program(
1656 context, &second_program, 1, &link_static_var_access);
1657 if (program == NULL || error != CL_SUCCESS)
1658 {
1659 log_error("ERROR: Unable to create a test program that tries to access "
1660 "a regular variable! (%s in %s:%d)\n",
1661 IGetErrorString(error), __FILE__, __LINE__);
1662 return -1;
1663 }
1664
1665 error = clCompileProgram(second_program, 1, &deviceID, NULL, 0, NULL, NULL,
1666 NULL, NULL);
1667 test_error(
1668 error,
1669 "Unable to compile a program that tries to access a regular variable");
1670
1671 cl_program two_programs[2] = { program, second_program };
1672 my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2,
1673 two_programs, NULL, NULL, &error);
1674 test_error(error,
1675 "clLinkProgram: Expected a different error code while linking a "
1676 "program that tries to access a regular variable");
1677
1678 /* All done! */
1679 error = clReleaseProgram(program);
1680 test_error(error, "Unable to release program object");
1681
1682 error = clReleaseProgram(second_program);
1683 test_error(error, "Unable to release program object");
1684
1685 error = clReleaseProgram(my_newly_linked_program);
1686 test_error(error, "Unable to release program object");
1687
1688 return 0;
1689 }
1690
test_two_file_regular_struct_access(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1691 int test_two_file_regular_struct_access(cl_device_id deviceID,
1692 cl_context context,
1693 cl_command_queue queue,
1694 int num_elements)
1695 {
1696 int error;
1697 cl_program program, second_program, my_newly_linked_program;
1698
1699 const char *sources[2] = {
1700 simple_kernel, compile_regular_struct
1701 }; // here we want to avoid linking error due to lack of kernels
1702 log_info("Compiling and linking two program objects, where one tries to "
1703 "access regular struct from another...\n");
1704 error = create_single_kernel_helper_create_program(context, &program, 2,
1705 sources);
1706 if (program == NULL || error != CL_SUCCESS)
1707 {
1708 log_error("ERROR: Unable to create a test program with regular struct! "
1709 "(%s in %s:%d)\n",
1710 IGetErrorString(error), __FILE__, __LINE__);
1711 return -1;
1712 }
1713
1714 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1715 NULL);
1716 test_error(error, "Unable to compile a simple program with regular struct");
1717
1718 error = create_single_kernel_helper_create_program(
1719 context, &second_program, 1, &link_static_struct_access);
1720 if (second_program == NULL || error != CL_SUCCESS)
1721 {
1722 log_error("ERROR: Unable to create a test program that tries to access "
1723 "a regular struct! (%s in %s:%d)\n",
1724 IGetErrorString(error), __FILE__, __LINE__);
1725 return -1;
1726 }
1727
1728 error = clCompileProgram(second_program, 1, &deviceID, NULL, 0, NULL, NULL,
1729 NULL, NULL);
1730 test_error(
1731 error,
1732 "Unable to compile a program that tries to access a regular struct");
1733
1734 cl_program two_programs[2] = { program, second_program };
1735 my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2,
1736 two_programs, NULL, NULL, &error);
1737 test_error(error,
1738 "clLinkProgram: Expected a different error code while linking a "
1739 "program that tries to access a regular struct");
1740
1741 /* All done! */
1742 error = clReleaseProgram(program);
1743 test_error(error, "Unable to release program object");
1744
1745 error = clReleaseProgram(second_program);
1746 test_error(error, "Unable to release program object");
1747
1748 error = clReleaseProgram(my_newly_linked_program);
1749 test_error(error, "Unable to release program object");
1750
1751 return 0;
1752 }
1753
1754
test_two_file_regular_function_access(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1755 int test_two_file_regular_function_access(cl_device_id deviceID,
1756 cl_context context,
1757 cl_command_queue queue,
1758 int num_elements)
1759 {
1760 int error;
1761 cl_program program, second_program, my_newly_linked_program;
1762
1763 const char *sources[2] = {
1764 simple_kernel, compile_regular_function
1765 }; // here we want to avoid linking error due to lack of kernels
1766 log_info("Compiling and linking two program objects, where one tries to "
1767 "access regular function from another...\n");
1768 error = create_single_kernel_helper_create_program(context, &program, 2,
1769 sources);
1770 if (program == NULL || error != CL_SUCCESS)
1771 {
1772 log_error("ERROR: Unable to create a test program with regular "
1773 "function! (%s in %s:%d)\n",
1774 IGetErrorString(error), __FILE__, __LINE__);
1775 return -1;
1776 }
1777
1778 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1779 NULL);
1780 test_error(error,
1781 "Unable to compile a simple program with regular function");
1782
1783 error = create_single_kernel_helper_create_program(
1784 context, &second_program, 1, &link_static_function_access);
1785 if (second_program == NULL || error != CL_SUCCESS)
1786 {
1787 log_error("ERROR: Unable to create a test program that tries to access "
1788 "a regular function! (%s in %s:%d)\n",
1789 IGetErrorString(error), __FILE__, __LINE__);
1790 return -1;
1791 }
1792
1793 error = clCompileProgram(second_program, 1, &deviceID, NULL, 0, NULL, NULL,
1794 NULL, NULL);
1795 test_error(
1796 error,
1797 "Unable to compile a program that tries to access a regular function");
1798
1799 cl_program two_programs[2] = { program, second_program };
1800 my_newly_linked_program = clLinkProgram(context, 1, &deviceID, NULL, 2,
1801 two_programs, NULL, NULL, &error);
1802 test_error(error,
1803 "clLinkProgram: Expected a different error code while linking a "
1804 "program that tries to access a regular function");
1805
1806 /* All done! */
1807 error = clReleaseProgram(program);
1808 test_error(error, "Unable to release program object");
1809
1810 error = clReleaseProgram(second_program);
1811 test_error(error, "Unable to release program object");
1812
1813 error = clReleaseProgram(my_newly_linked_program);
1814 test_error(error, "Unable to release program object");
1815
1816 return 0;
1817 }
1818
test_simple_embedded_header_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1819 int test_simple_embedded_header_link(cl_device_id deviceID, cl_context context,
1820 cl_command_queue queue, int num_elements)
1821 {
1822 int error;
1823 cl_program program, header, simple_program;
1824
1825 log_info("Testing a simple embedded header link...\n");
1826 program = clCreateProgramWithSource(
1827 context, 1, &another_simple_kernel_with_header, NULL, &error);
1828 if (program == NULL || error != CL_SUCCESS)
1829 {
1830 log_error(
1831 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
1832 IGetErrorString(error), __FILE__, __LINE__);
1833 return -1;
1834 }
1835
1836 header =
1837 clCreateProgramWithSource(context, 1, &simple_header, NULL, &error);
1838 if (header == NULL || error != CL_SUCCESS)
1839 {
1840 log_error(
1841 "ERROR: Unable to create a simple header program! (%s in %s:%d)\n",
1842 IGetErrorString(error), __FILE__, __LINE__);
1843 return -1;
1844 }
1845
1846 error = clCompileProgram(program, 1, &deviceID, NULL, 1, &header,
1847 &simple_header_name, NULL, NULL);
1848 test_error(error,
1849 "Unable to compile a simple program with embedded header");
1850
1851 error = create_single_kernel_helper_create_program(context, &simple_program,
1852 1, &simple_kernel);
1853 if (simple_program == NULL || error != CL_SUCCESS)
1854 {
1855 log_error(
1856 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
1857 IGetErrorString(error), __FILE__, __LINE__);
1858 return -1;
1859 }
1860
1861 error = clCompileProgram(simple_program, 1, &deviceID, NULL, 0, NULL, NULL,
1862 NULL, NULL);
1863 test_error(error, "Unable to compile a simple program");
1864
1865 cl_program two_programs[2] = { program, simple_program };
1866 cl_program fully_linked_program = clLinkProgram(
1867 context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
1868 test_error(error,
1869 "Unable to create an executable from two binaries, one compiled "
1870 "with embedded header");
1871
1872 /* All done! */
1873 error = clReleaseProgram(program);
1874 test_error(error, "Unable to release program object");
1875
1876 error = clReleaseProgram(header);
1877 test_error(error, "Unable to release program object");
1878
1879 error = clReleaseProgram(simple_program);
1880 test_error(error, "Unable to release program object");
1881
1882 error = clReleaseProgram(fully_linked_program);
1883 test_error(error, "Unable to release program object");
1884
1885 return 0;
1886 }
1887
1888 const char *when_i_pondered_weak_and_weary = "When I pondered weak and weary!";
1889
simple_link_callback(cl_program program,void * user_data)1890 static void CL_CALLBACK simple_link_callback(cl_program program,
1891 void *user_data)
1892 {
1893 simple_user_data *simple_link_user_data = (simple_user_data *)user_data;
1894 log_info("in the simple_link_callback: program %p just completed linking "
1895 "with '%s'\n",
1896 program, (const char *)simple_link_user_data->m_message);
1897 if (strcmp(when_i_pondered_weak_and_weary, simple_link_user_data->m_message)
1898 != 0)
1899 {
1900 log_error("ERROR: in the simple_compile_callback: Expected '%s' and "
1901 "got %s! (in %s:%d)\n",
1902 when_i_pondered_weak_and_weary,
1903 simple_link_user_data->m_message, __FILE__, __LINE__);
1904 }
1905
1906 int error;
1907 log_info("in the simple_link_callback: program %p just completed linking "
1908 "with '%p'\n",
1909 program, simple_link_user_data->m_event);
1910
1911 error = clSetUserEventStatus(simple_link_user_data->m_event, CL_COMPLETE);
1912 if (error != CL_SUCCESS)
1913 {
1914 log_error("ERROR: simple_link_callback: Unable to set user event "
1915 "status to CL_COMPLETE! (%s in %s:%d)\n",
1916 IGetErrorString(error), __FILE__, __LINE__);
1917 exit(-1);
1918 }
1919 log_info("in the simple_link_callback: Successfully signaled "
1920 "link_program_completion_event event!\n");
1921 }
1922
test_simple_link_with_callback(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)1923 int test_simple_link_with_callback(cl_device_id deviceID, cl_context context,
1924 cl_command_queue queue, int num_elements)
1925 {
1926 int error;
1927 cl_program program;
1928 cl_event link_program_completion_event;
1929
1930 log_info("Testing a simple linking with callback...\n");
1931 error = create_single_kernel_helper_create_program(context, &program, 1,
1932 &simple_kernel);
1933 if (program == NULL || error != CL_SUCCESS)
1934 {
1935 log_error(
1936 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
1937 IGetErrorString(error), __FILE__, __LINE__);
1938 return -1;
1939 }
1940
1941 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
1942 NULL);
1943 test_error(error, "Unable to compile a simple program");
1944
1945 link_program_completion_event = clCreateUserEvent(context, &error);
1946 test_error(error, "Unable to create a user event");
1947
1948 simple_user_data simple_link_user_data = { when_i_pondered_weak_and_weary,
1949 link_program_completion_event };
1950
1951 cl_program my_linked_library = clLinkProgram(
1952 context, 1, &deviceID, NULL, 1, &program, simple_link_callback,
1953 (void *)&simple_link_user_data, &error);
1954 test_error(error, "Unable to link a simple program");
1955
1956 error = clWaitForEvents(1, &link_program_completion_event);
1957 test_error(
1958 error,
1959 "clWaitForEvents failed when waiting on link_program_completion_event");
1960
1961 /* All done! */
1962 error = clReleaseEvent(link_program_completion_event);
1963 test_error(error, "Unable to release event object");
1964
1965 error = clReleaseProgram(program);
1966 test_error(error, "Unable to release program object");
1967
1968 error = clReleaseProgram(my_linked_library);
1969 test_error(error, "Unable to release program object");
1970
1971 return 0;
1972 }
1973
initBuffer(float * & srcBuffer,unsigned int cnDimension)1974 static void initBuffer(float *&srcBuffer, unsigned int cnDimension)
1975 {
1976 float num = 0.0f;
1977
1978 for (unsigned int i = 0; i < cnDimension; i++)
1979 {
1980 if ((i % 10) == 0)
1981 {
1982 num = 0.0f;
1983 }
1984
1985 srcBuffer[i] = num;
1986 num = num + 1.0f;
1987 }
1988 }
1989
verifyCopyBuffer(cl_context context,cl_command_queue queue,cl_kernel kernel)1990 static int verifyCopyBuffer(cl_context context, cl_command_queue queue,
1991 cl_kernel kernel)
1992 {
1993 int error, result = CL_SUCCESS;
1994 const size_t cnDimension = 32;
1995
1996 // Allocate source buffer
1997 float *srcBuffer = (float *)malloc(cnDimension * sizeof(float));
1998 float *dstBuffer = (float *)malloc(cnDimension * sizeof(float));
1999
2000 if (srcBuffer == NULL)
2001 {
2002 log_error("ERROR: Unable to allocate srcBuffer float array with %lu "
2003 "floats! (in %s:%d)\n",
2004 cnDimension, __FILE__, __LINE__);
2005 return -1;
2006 }
2007 if (dstBuffer == NULL)
2008 {
2009 log_error("ERROR: Unable to allocate dstBuffer float array with %lu "
2010 "floats! (in %s:%d)\n",
2011 cnDimension, __FILE__, __LINE__);
2012 return -1;
2013 }
2014
2015 if (srcBuffer && dstBuffer)
2016 {
2017 // initialize host memory
2018 initBuffer(srcBuffer, cnDimension);
2019
2020 // Allocate device memory
2021 cl_mem deviceMemSrc =
2022 clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
2023 cnDimension * sizeof(cl_float), srcBuffer, &error);
2024 test_error(error, "Unable to create a source memory buffer");
2025
2026 cl_mem deviceMemDst =
2027 clCreateBuffer(context, CL_MEM_WRITE_ONLY,
2028 cnDimension * sizeof(cl_float), 0, &error);
2029 test_error(error, "Unable to create a destination memory buffer");
2030
2031 // Set kernel args
2032 // Set parameter 0 to be the source buffer
2033 error =
2034 clSetKernelArg(kernel, 0, sizeof(cl_mem), (void *)&deviceMemSrc);
2035 test_error(error, "Unable to set the first kernel argument");
2036
2037 // Set parameter 1 to be the destination buffer
2038 error =
2039 clSetKernelArg(kernel, 1, sizeof(cl_mem), (void *)&deviceMemDst);
2040 test_error(error, "Unable to set the second kernel argument");
2041
2042 // Execute kernel
2043 error = clEnqueueNDRangeKernel(queue, kernel, 1, NULL, &cnDimension, 0,
2044 0, NULL, NULL);
2045 test_error(error, "Unable to enqueue kernel");
2046
2047 error = clFlush(queue);
2048 test_error(error, "Unable to flush the queue");
2049
2050 // copy results from device back to host
2051 error = clEnqueueReadBuffer(queue, deviceMemDst, CL_TRUE, 0,
2052 cnDimension * sizeof(cl_float), dstBuffer,
2053 0, NULL, NULL);
2054 test_error(error, "Unable to read the destination buffer");
2055
2056 error = clFlush(queue);
2057 test_error(error, "Unable to flush the queue");
2058
2059 // Compare the source and destination buffers
2060 const int *pSrc = (int *)srcBuffer;
2061 const int *pDst = (int *)dstBuffer;
2062 int mismatch = 0;
2063
2064 for (size_t i = 0; i < cnDimension; i++)
2065 {
2066 if (pSrc[i] != pDst[i])
2067 {
2068 if (mismatch < 4)
2069 {
2070 log_info("Offset %08lX: Expected %08X, Got %08X\n", i * 4,
2071 pSrc[i], pDst[i]);
2072 }
2073 else
2074 {
2075 log_info(".");
2076 }
2077 mismatch++;
2078 }
2079 }
2080
2081 if (mismatch)
2082 {
2083 log_info("*** %d mismatches found, TEST FAILS! ***\n", mismatch);
2084 result = -1;
2085 }
2086 else
2087 {
2088 log_info("Buffers match, test passes.\n");
2089 }
2090
2091 free(srcBuffer);
2092 srcBuffer = NULL;
2093 free(dstBuffer);
2094 dstBuffer = NULL;
2095
2096 if (deviceMemSrc)
2097 {
2098 error = clReleaseMemObject(deviceMemSrc);
2099 test_error(error, "Unable to release memory object");
2100 }
2101
2102 if (deviceMemDst)
2103 {
2104 error = clReleaseMemObject(deviceMemDst);
2105 test_error(error, "Unable to release memory object");
2106 }
2107 }
2108 return result;
2109 }
2110
test_execute_after_simple_compile_and_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2111 int test_execute_after_simple_compile_and_link(cl_device_id deviceID,
2112 cl_context context,
2113 cl_command_queue queue,
2114 int num_elements)
2115 {
2116 int error;
2117 cl_program program;
2118
2119 log_info("Testing execution after a simple compile and link...\n");
2120 error = create_single_kernel_helper_create_program(context, &program, 1,
2121 &simple_kernel);
2122 if (program == NULL || error != CL_SUCCESS)
2123 {
2124 log_error(
2125 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2126 IGetErrorString(error), __FILE__, __LINE__);
2127 return -1;
2128 }
2129
2130 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2131 NULL);
2132 test_error(error, "Unable to compile a simple program");
2133
2134 cl_program my_newly_linked_program = clLinkProgram(
2135 context, 1, &deviceID, NULL, 1, &program, NULL, NULL, &error);
2136 test_error(error, "Unable to link a simple program");
2137
2138 cl_kernel kernel =
2139 clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
2140 test_error(error, "Unable to create a simple kernel");
2141
2142 error = verifyCopyBuffer(context, queue, kernel);
2143 if (error != CL_SUCCESS) return error;
2144
2145 /* All done! */
2146 error = clReleaseKernel(kernel);
2147 test_error(error, "Unable to release kernel object");
2148
2149 error = clReleaseProgram(program);
2150 test_error(error, "Unable to release program object");
2151
2152 error = clReleaseProgram(my_newly_linked_program);
2153 test_error(error, "Unable to release program object");
2154
2155 return 0;
2156 }
2157
test_execute_after_simple_compile_and_link_no_device_info(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2158 int test_execute_after_simple_compile_and_link_no_device_info(
2159 cl_device_id deviceID, cl_context context, cl_command_queue queue,
2160 int num_elements)
2161 {
2162 int error;
2163 cl_program program;
2164
2165 log_info("Testing execution after a simple compile and link with no device "
2166 "information provided...\n");
2167 error = create_single_kernel_helper_create_program(context, &program, 1,
2168 &simple_kernel);
2169 if (program == NULL || error != CL_SUCCESS)
2170 {
2171 log_error(
2172 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2173 IGetErrorString(error), __FILE__, __LINE__);
2174 return -1;
2175 }
2176
2177 error = clCompileProgram(program, 0, NULL, NULL, 0, NULL, NULL, NULL, NULL);
2178 test_error(error, "Unable to compile a simple program");
2179
2180 cl_program my_newly_linked_program =
2181 clLinkProgram(context, 0, NULL, NULL, 1, &program, NULL, NULL, &error);
2182 test_error(error, "Unable to link a simple program");
2183
2184 cl_kernel kernel =
2185 clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
2186 test_error(error, "Unable to create a simple kernel");
2187
2188 error = verifyCopyBuffer(context, queue, kernel);
2189 if (error != CL_SUCCESS) return error;
2190
2191 /* All done! */
2192 error = clReleaseKernel(kernel);
2193 test_error(error, "Unable to release kernel object");
2194
2195 error = clReleaseProgram(program);
2196 test_error(error, "Unable to release program object");
2197
2198 error = clReleaseProgram(my_newly_linked_program);
2199 test_error(error, "Unable to release program object");
2200
2201 return 0;
2202 }
2203
test_execute_after_simple_compile_and_link_with_defines(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2204 int test_execute_after_simple_compile_and_link_with_defines(
2205 cl_device_id deviceID, cl_context context, cl_command_queue queue,
2206 int num_elements)
2207 {
2208 int error;
2209 cl_program program;
2210
2211 log_info(
2212 "Testing execution after a simple compile and link with defines...\n");
2213 error = create_single_kernel_helper_create_program(
2214 context, &program, 1, &simple_kernel_with_defines,
2215 "-DFIRST=5 -DSECOND=37");
2216 if (program == NULL || error != CL_SUCCESS)
2217 {
2218 log_error(
2219 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2220 IGetErrorString(error), __FILE__, __LINE__);
2221 return -1;
2222 }
2223
2224 error = clCompileProgram(program, 1, &deviceID, "-DFIRST=5 -DSECOND=37", 0,
2225 NULL, NULL, NULL, NULL);
2226 test_error(error, "Unable to compile a simple program");
2227
2228 cl_program my_newly_linked_program = clLinkProgram(
2229 context, 1, &deviceID, NULL, 1, &program, NULL, NULL, &error);
2230 test_error(error, "Unable to link a simple program");
2231
2232 cl_kernel kernel =
2233 clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
2234 test_error(error, "Unable to create a simple kernel");
2235
2236 error = verifyCopyBuffer(context, queue, kernel);
2237 if (error != CL_SUCCESS) return error;
2238
2239 /* All done! */
2240 error = clReleaseKernel(kernel);
2241 test_error(error, "Unable to release kernel object");
2242
2243 error = clReleaseProgram(program);
2244 test_error(error, "Unable to release program object");
2245
2246 error = clReleaseProgram(my_newly_linked_program);
2247 test_error(error, "Unable to release program object");
2248
2249 return 0;
2250 }
2251
test_execute_after_serialize_reload_object(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2252 int test_execute_after_serialize_reload_object(cl_device_id deviceID,
2253 cl_context context,
2254 cl_command_queue queue,
2255 int num_elements)
2256 {
2257 int error;
2258 cl_program program;
2259 size_t binarySize;
2260 unsigned char *binary;
2261
2262 log_info("Testing execution after serialization and reloading of the "
2263 "object...\n");
2264 error = create_single_kernel_helper_create_program(context, &program, 1,
2265 &simple_kernel);
2266 if (program == NULL || error != CL_SUCCESS)
2267 {
2268 log_error(
2269 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2270 IGetErrorString(error), __FILE__, __LINE__);
2271 return -1;
2272 }
2273
2274 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2275 NULL);
2276 test_error(error, "Unable to compile a simple program");
2277
2278 // Get the size of the resulting binary (only one device)
2279 error = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES,
2280 sizeof(binarySize), &binarySize, NULL);
2281 test_error(error, "Unable to get binary size");
2282
2283 // Sanity check
2284 if (binarySize == 0)
2285 {
2286 log_error("ERROR: Binary size of program is zero (in %s:%d)\n",
2287 __FILE__, __LINE__);
2288 return -1;
2289 }
2290
2291 // Create a buffer and get the actual binary
2292 binary = (unsigned char *)malloc(sizeof(unsigned char) * binarySize);
2293 if (binary == NULL)
2294 {
2295 log_error("ERROR: Unable to allocate binary character array with %lu "
2296 "characters! (in %s:%d)\n",
2297 binarySize, __FILE__, __LINE__);
2298 return -1;
2299 }
2300
2301 unsigned char *buffers[1] = { binary };
2302 cl_int loadErrors[1];
2303
2304 // Do another sanity check here first
2305 size_t size;
2306 error = clGetProgramInfo(program, CL_PROGRAM_BINARIES, 0, NULL, &size);
2307 test_error(error, "Unable to get expected size of binaries array");
2308 if (size != sizeof(buffers))
2309 {
2310 log_error("ERROR: Expected size of binaries array in clGetProgramInfo "
2311 "is incorrect (should be %d, got %d) (in %s:%d)\n",
2312 (int)sizeof(buffers), (int)size, __FILE__, __LINE__);
2313 free(binary);
2314 return -1;
2315 }
2316
2317 error = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(buffers),
2318 &buffers, NULL);
2319 test_error(error, "Unable to get program binary");
2320
2321 // use clCreateProgramWithBinary
2322 cl_program program_with_binary = clCreateProgramWithBinary(
2323 context, 1, &deviceID, &binarySize, (const unsigned char **)buffers,
2324 loadErrors, &error);
2325 test_error(error, "Unable to create program with binary");
2326
2327 cl_program my_newly_linked_program =
2328 clLinkProgram(context, 1, &deviceID, NULL, 1, &program_with_binary,
2329 NULL, NULL, &error);
2330 test_error(error, "Unable to link a simple program");
2331
2332 cl_kernel kernel =
2333 clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
2334 test_error(error, "Unable to create a simple kernel");
2335
2336 error = verifyCopyBuffer(context, queue, kernel);
2337 if (error != CL_SUCCESS) return error;
2338
2339 /* All done! */
2340 error = clReleaseKernel(kernel);
2341 test_error(error, "Unable to release kernel object");
2342
2343 error = clReleaseProgram(program);
2344 test_error(error, "Unable to release program object");
2345
2346 error = clReleaseProgram(my_newly_linked_program);
2347 test_error(error, "Unable to release program object");
2348
2349 error = clReleaseProgram(program_with_binary);
2350 test_error(error, "Unable to release program object");
2351
2352 free(binary);
2353
2354 return 0;
2355 }
2356
test_execute_after_serialize_reload_library(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2357 int test_execute_after_serialize_reload_library(cl_device_id deviceID,
2358 cl_context context,
2359 cl_command_queue queue,
2360 int num_elements)
2361 {
2362 int error;
2363 cl_program program, another_program;
2364 size_t binarySize;
2365 unsigned char *binary;
2366
2367 log_info(
2368 "Testing execution after linking a binary with a simple library...\n");
2369 // we will test creation of a simple library from one file
2370 error = create_single_kernel_helper_create_program(context, &program, 1,
2371 &simple_kernel);
2372 if (program == NULL || error != CL_SUCCESS)
2373 {
2374 log_error(
2375 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2376 IGetErrorString(error), __FILE__, __LINE__);
2377 return -1;
2378 }
2379
2380 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2381 NULL);
2382 test_error(error, "Unable to compile a simple program");
2383
2384 cl_program my_newly_minted_library =
2385 clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program,
2386 NULL, NULL, &error);
2387 test_error(error, "Unable to create a simple library");
2388
2389
2390 // Get the size of the resulting library (only one device)
2391 error = clGetProgramInfo(my_newly_minted_library, CL_PROGRAM_BINARY_SIZES,
2392 sizeof(binarySize), &binarySize, NULL);
2393 test_error(error, "Unable to get binary size");
2394
2395 // Sanity check
2396 if (binarySize == 0)
2397 {
2398 log_error("ERROR: Binary size of program is zero (in %s:%d)\n",
2399 __FILE__, __LINE__);
2400 return -1;
2401 }
2402
2403 // Create a buffer and get the actual binary
2404 binary = (unsigned char *)malloc(sizeof(unsigned char) * binarySize);
2405 if (binary == NULL)
2406 {
2407 log_error("ERROR: Unable to allocate binary character array with %lu "
2408 "characters (in %s:%d)!",
2409 binarySize, __FILE__, __LINE__);
2410 return -1;
2411 }
2412 unsigned char *buffers[1] = { binary };
2413 cl_int loadErrors[1];
2414
2415 // Do another sanity check here first
2416 size_t size;
2417 error = clGetProgramInfo(my_newly_minted_library, CL_PROGRAM_BINARIES, 0,
2418 NULL, &size);
2419 test_error(error, "Unable to get expected size of binaries array");
2420 if (size != sizeof(buffers))
2421 {
2422 log_error("ERROR: Expected size of binaries array in clGetProgramInfo "
2423 "is incorrect (should be %d, got %d) (in %s:%d)\n",
2424 (int)sizeof(buffers), (int)size, __FILE__, __LINE__);
2425 free(binary);
2426 return -1;
2427 }
2428
2429 error = clGetProgramInfo(my_newly_minted_library, CL_PROGRAM_BINARIES,
2430 sizeof(buffers), &buffers, NULL);
2431 test_error(error, "Unable to get program binary");
2432
2433 // use clCreateProgramWithBinary
2434 cl_program library_with_binary = clCreateProgramWithBinary(
2435 context, 1, &deviceID, &binarySize, (const unsigned char **)buffers,
2436 loadErrors, &error);
2437 test_error(error, "Unable to create program with binary");
2438
2439 error = create_single_kernel_helper_create_program(
2440 context, &another_program, 1, &another_simple_kernel);
2441 if (another_program == NULL || error != CL_SUCCESS)
2442 {
2443 log_error(
2444 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2445 IGetErrorString(error), __FILE__, __LINE__);
2446 return -1;
2447 }
2448
2449 error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL,
2450 NULL, NULL);
2451 test_error(error, "Unable to compile a simple program");
2452
2453 cl_program program_and_archive[2] = { another_program,
2454 library_with_binary };
2455 cl_program fully_linked_program = clLinkProgram(
2456 context, 1, &deviceID, "", 2, program_and_archive, NULL, NULL, &error);
2457 test_error(error,
2458 "Unable to create an executable from a binary and a library");
2459
2460 cl_kernel kernel =
2461 clCreateKernel(fully_linked_program, "CopyBuffer", &error);
2462 test_error(error, "Unable to create a simple kernel");
2463
2464 error = verifyCopyBuffer(context, queue, kernel);
2465 if (error != CL_SUCCESS) return error;
2466
2467 cl_kernel another_kernel =
2468 clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
2469 test_error(error, "Unable to create another simple kernel");
2470
2471 error = verifyCopyBuffer(context, queue, another_kernel);
2472 if (error != CL_SUCCESS) return error;
2473
2474 /* All done! */
2475 error = clReleaseKernel(kernel);
2476 test_error(error, "Unable to release kernel object");
2477
2478 error = clReleaseKernel(another_kernel);
2479 test_error(error, "Unable to release another kernel object");
2480
2481 error = clReleaseProgram(program);
2482 test_error(error, "Unable to release program object");
2483
2484 error = clReleaseProgram(another_program);
2485 test_error(error, "Unable to release program object");
2486
2487 error = clReleaseProgram(my_newly_minted_library);
2488 test_error(error, "Unable to release program object");
2489
2490 error = clReleaseProgram(library_with_binary);
2491 test_error(error, "Unable to release program object");
2492
2493 error = clReleaseProgram(fully_linked_program);
2494 test_error(error, "Unable to release program object");
2495
2496 free(binary);
2497
2498 return 0;
2499 }
2500
program_compile_completion_callback(cl_program program,void * user_data)2501 static void CL_CALLBACK program_compile_completion_callback(cl_program program,
2502 void *user_data)
2503 {
2504 int error;
2505 cl_event compile_program_completion_event = (cl_event)user_data;
2506 log_info("in the program_compile_completion_callback: program %p just "
2507 "completed compiling with '%p'\n",
2508 program, compile_program_completion_event);
2509
2510 error = clSetUserEventStatus(compile_program_completion_event, CL_COMPLETE);
2511 if (error != CL_SUCCESS)
2512 {
2513 log_error("ERROR: in the program_compile_completion_callback: Unable "
2514 "to set user event status to CL_COMPLETE! (%s in %s:%d)\n",
2515 IGetErrorString(error), __FILE__, __LINE__);
2516 exit(-1);
2517 }
2518 log_info("in the program_compile_completion_callback: Successfully "
2519 "signaled compile_program_completion_event event!\n");
2520 }
2521
program_link_completion_callback(cl_program program,void * user_data)2522 static void CL_CALLBACK program_link_completion_callback(cl_program program,
2523 void *user_data)
2524 {
2525 int error;
2526 cl_event link_program_completion_event = (cl_event)user_data;
2527 log_info("in the program_link_completion_callback: program %p just "
2528 "completed linking with '%p'\n",
2529 program, link_program_completion_event);
2530
2531 error = clSetUserEventStatus(link_program_completion_event, CL_COMPLETE);
2532 if (error != CL_SUCCESS)
2533 {
2534 log_error("ERROR: in the program_link_completion_callback: Unable to "
2535 "set user event status to CL_COMPLETE! (%s in %s:%d)\n",
2536 IGetErrorString(error), __FILE__, __LINE__);
2537 exit(-1);
2538 }
2539 log_info("in the program_link_completion_callback: Successfully signaled "
2540 "link_program_completion_event event!\n");
2541 }
2542
test_execute_after_simple_compile_and_link_with_callbacks(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2543 int test_execute_after_simple_compile_and_link_with_callbacks(
2544 cl_device_id deviceID, cl_context context, cl_command_queue queue,
2545 int num_elements)
2546 {
2547 int error;
2548 cl_program program;
2549 cl_event compile_program_completion_event, link_program_completion_event;
2550
2551 log_info("Testing execution after a simple compile and link with "
2552 "callbacks...\n");
2553 error = create_single_kernel_helper_create_program(context, &program, 1,
2554 &simple_kernel);
2555 if (program == NULL || error != CL_SUCCESS)
2556 {
2557 log_error(
2558 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2559 IGetErrorString(error), __FILE__, __LINE__);
2560 return -1;
2561 }
2562
2563 compile_program_completion_event = clCreateUserEvent(context, &error);
2564 test_error(error, "Unable to create a user event");
2565
2566 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL,
2567 program_compile_completion_callback,
2568 (void *)compile_program_completion_event);
2569 test_error(error, "Unable to compile a simple program");
2570
2571 error = clWaitForEvents(1, &compile_program_completion_event);
2572 test_error(error,
2573 "clWaitForEvents failed when waiting on "
2574 "compile_program_completion_event");
2575
2576 error = clReleaseEvent(compile_program_completion_event);
2577 test_error(error, "Unable to release event object");
2578
2579 link_program_completion_event = clCreateUserEvent(context, &error);
2580 test_error(error, "Unable to create a user event");
2581
2582 cl_program my_newly_linked_program =
2583 clLinkProgram(context, 1, &deviceID, NULL, 1, &program,
2584 program_link_completion_callback,
2585 (void *)link_program_completion_event, &error);
2586 test_error(error, "Unable to link a simple program");
2587
2588 error = clWaitForEvents(1, &link_program_completion_event);
2589 test_error(
2590 error,
2591 "clWaitForEvents failed when waiting on link_program_completion_event");
2592
2593 error = clReleaseEvent(link_program_completion_event);
2594 test_error(error, "Unable to release event object");
2595
2596 cl_kernel kernel =
2597 clCreateKernel(my_newly_linked_program, "CopyBuffer", &error);
2598 test_error(error, "Unable to create a simple kernel");
2599
2600 error = verifyCopyBuffer(context, queue, kernel);
2601 if (error != CL_SUCCESS) return error;
2602
2603 /* All done! */
2604 error = clReleaseKernel(kernel);
2605 test_error(error, "Unable to release kernel object");
2606
2607 error = clReleaseProgram(program);
2608 test_error(error, "Unable to release program object");
2609
2610 error = clReleaseProgram(my_newly_linked_program);
2611 test_error(error, "Unable to release program object");
2612
2613 return 0;
2614 }
2615
test_simple_library_only(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2616 int test_simple_library_only(cl_device_id deviceID, cl_context context,
2617 cl_command_queue queue, int num_elements)
2618 {
2619 int error;
2620 cl_program program;
2621
2622 log_info("Testing creation of a simple library...\n");
2623 error = create_single_kernel_helper_create_program(context, &program, 1,
2624 &simple_kernel);
2625 if (program == NULL || error != CL_SUCCESS)
2626 {
2627 log_error(
2628 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2629 IGetErrorString(error), __FILE__, __LINE__);
2630 return -1;
2631 }
2632
2633 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2634 NULL);
2635 test_error(error, "Unable to compile a simple program");
2636
2637 cl_program my_newly_minted_library =
2638 clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program,
2639 NULL, NULL, &error);
2640 test_error(error, "Unable to create a simple library");
2641
2642 /* All done! */
2643 error = clReleaseProgram(program);
2644 test_error(error, "Unable to release program object");
2645
2646 error = clReleaseProgram(my_newly_minted_library);
2647 test_error(error, "Unable to release program object");
2648
2649 return 0;
2650 }
2651
test_simple_library_with_callback(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2652 int test_simple_library_with_callback(cl_device_id deviceID, cl_context context,
2653 cl_command_queue queue, int num_elements)
2654 {
2655 int error;
2656 cl_program program;
2657 cl_event link_program_completion_event;
2658
2659 log_info("Testing creation of a simple library with a callback...\n");
2660 error = create_single_kernel_helper_create_program(context, &program, 1,
2661 &simple_kernel);
2662 if (program == NULL || error != CL_SUCCESS)
2663 {
2664 log_error(
2665 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2666 IGetErrorString(error), __FILE__, __LINE__);
2667 return -1;
2668 }
2669
2670 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2671 NULL);
2672 test_error(error, "Unable to compile a simple program");
2673
2674 link_program_completion_event = clCreateUserEvent(context, &error);
2675 test_error(error, "Unable to create a user event");
2676
2677 simple_user_data simple_link_user_data = { when_i_pondered_weak_and_weary,
2678 link_program_completion_event };
2679
2680 cl_program my_newly_minted_library = clLinkProgram(
2681 context, 1, &deviceID, "-create-library", 1, &program,
2682 simple_link_callback, (void *)&simple_link_user_data, &error);
2683 test_error(error, "Unable to create a simple library");
2684
2685 error = clWaitForEvents(1, &link_program_completion_event);
2686 test_error(
2687 error,
2688 "clWaitForEvents failed when waiting on link_program_completion_event");
2689
2690 /* All done! */
2691 error = clReleaseEvent(link_program_completion_event);
2692 test_error(error, "Unable to release event object");
2693
2694 error = clReleaseProgram(program);
2695 test_error(error, "Unable to release program object");
2696
2697 error = clReleaseProgram(my_newly_minted_library);
2698 test_error(error, "Unable to release program object");
2699
2700 return 0;
2701 }
2702
test_simple_library_with_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2703 int test_simple_library_with_link(cl_device_id deviceID, cl_context context,
2704 cl_command_queue queue, int num_elements)
2705 {
2706 int error;
2707 cl_program program, another_program;
2708
2709 log_info("Testing creation and linking with a simple library...\n");
2710 error = create_single_kernel_helper_create_program(context, &program, 1,
2711 &simple_kernel);
2712 if (program == NULL || error != CL_SUCCESS)
2713 {
2714 log_error(
2715 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2716 IGetErrorString(error), __FILE__, __LINE__);
2717 return -1;
2718 }
2719
2720 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2721 NULL);
2722 test_error(error, "Unable to compile a simple program");
2723
2724 cl_program my_newly_minted_library =
2725 clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program,
2726 NULL, NULL, &error);
2727 test_error(error, "Unable to create a simple library");
2728
2729 error = create_single_kernel_helper_create_program(
2730 context, &another_program, 1, &another_simple_kernel);
2731 if (another_program == NULL || error != CL_SUCCESS)
2732 {
2733 log_error(
2734 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2735 IGetErrorString(error), __FILE__, __LINE__);
2736 return -1;
2737 }
2738
2739 error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL,
2740 NULL, NULL);
2741 test_error(error, "Unable to compile a simple program");
2742
2743 cl_program program_and_archive[2] = { another_program,
2744 my_newly_minted_library };
2745 cl_program fully_linked_program = clLinkProgram(
2746 context, 1, &deviceID, "", 2, program_and_archive, NULL, NULL, &error);
2747 test_error(error,
2748 "Unable to create an executable from a binary and a library");
2749
2750 /* All done! */
2751 error = clReleaseProgram(program);
2752 test_error(error, "Unable to release program object");
2753
2754 error = clReleaseProgram(another_program);
2755 test_error(error, "Unable to release program object");
2756
2757 error = clReleaseProgram(my_newly_minted_library);
2758 test_error(error, "Unable to release program object");
2759
2760 error = clReleaseProgram(fully_linked_program);
2761 test_error(error, "Unable to release program object");
2762
2763 return 0;
2764 }
2765
test_execute_after_simple_library_with_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2766 int test_execute_after_simple_library_with_link(cl_device_id deviceID,
2767 cl_context context,
2768 cl_command_queue queue,
2769 int num_elements)
2770 {
2771 int error;
2772 cl_program program, another_program;
2773
2774 log_info(
2775 "Testing execution after linking a binary with a simple library...\n");
2776 error = create_single_kernel_helper_create_program(context, &program, 1,
2777 &simple_kernel);
2778 if (program == NULL || error != CL_SUCCESS)
2779 {
2780 log_error(
2781 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2782 IGetErrorString(error), __FILE__, __LINE__);
2783 return -1;
2784 }
2785
2786 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2787 NULL);
2788 test_error(error, "Unable to compile a simple program");
2789
2790 cl_program my_newly_minted_library =
2791 clLinkProgram(context, 1, &deviceID, "-create-library", 1, &program,
2792 NULL, NULL, &error);
2793 test_error(error, "Unable to create a simple library");
2794
2795 error = create_single_kernel_helper_create_program(
2796 context, &another_program, 1, &another_simple_kernel);
2797 if (another_program == NULL || error != CL_SUCCESS)
2798 {
2799 log_error(
2800 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2801 IGetErrorString(error), __FILE__, __LINE__);
2802 return -1;
2803 }
2804
2805 error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL,
2806 NULL, NULL);
2807 test_error(error, "Unable to compile a simple program");
2808
2809 cl_program program_and_archive[2] = { another_program,
2810 my_newly_minted_library };
2811 cl_program fully_linked_program = clLinkProgram(
2812 context, 1, &deviceID, "", 2, program_and_archive, NULL, NULL, &error);
2813 test_error(error,
2814 "Unable to create an executable from a binary and a library");
2815
2816 cl_kernel kernel =
2817 clCreateKernel(fully_linked_program, "CopyBuffer", &error);
2818 test_error(error, "Unable to create a simple kernel");
2819
2820 error = verifyCopyBuffer(context, queue, kernel);
2821 if (error != CL_SUCCESS) return error;
2822
2823 cl_kernel another_kernel =
2824 clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
2825 test_error(error, "Unable to create another simple kernel");
2826
2827 error = verifyCopyBuffer(context, queue, another_kernel);
2828 if (error != CL_SUCCESS) return error;
2829
2830 /* All done! */
2831 error = clReleaseKernel(kernel);
2832 test_error(error, "Unable to release kernel object");
2833
2834 error = clReleaseKernel(another_kernel);
2835 test_error(error, "Unable to release another kernel object");
2836
2837 error = clReleaseProgram(program);
2838 test_error(error, "Unable to release program object");
2839
2840 error = clReleaseProgram(another_program);
2841 test_error(error, "Unable to release program object");
2842
2843 error = clReleaseProgram(my_newly_minted_library);
2844 test_error(error, "Unable to release program object");
2845
2846 error = clReleaseProgram(fully_linked_program);
2847 test_error(error, "Unable to release program object");
2848
2849 return 0;
2850 }
2851
test_two_file_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2852 int test_two_file_link(cl_device_id deviceID, cl_context context,
2853 cl_command_queue queue, int num_elements)
2854 {
2855 int error;
2856 cl_program program, another_program;
2857
2858 log_info("Testing two file compiling and linking...\n");
2859 error = create_single_kernel_helper_create_program(context, &program, 1,
2860 &simple_kernel);
2861 if (program == NULL || error != CL_SUCCESS)
2862 {
2863 log_error(
2864 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2865 IGetErrorString(error), __FILE__, __LINE__);
2866 return -1;
2867 }
2868
2869 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2870 NULL);
2871 test_error(error, "Unable to compile a simple program");
2872
2873
2874 error = create_single_kernel_helper_create_program(
2875 context, &another_program, 1, &another_simple_kernel);
2876 if (another_program == NULL || error != CL_SUCCESS)
2877 {
2878 log_error(
2879 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2880 IGetErrorString(error), __FILE__, __LINE__);
2881 return -1;
2882 }
2883
2884 error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL,
2885 NULL, NULL);
2886 test_error(error, "Unable to compile a simple program");
2887
2888 cl_program two_programs[2] = { program, another_program };
2889 cl_program fully_linked_program = clLinkProgram(
2890 context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
2891 test_error(error, "Unable to create an executable from two binaries");
2892
2893 /* All done! */
2894 error = clReleaseProgram(program);
2895 test_error(error, "Unable to release program object");
2896
2897 error = clReleaseProgram(another_program);
2898 test_error(error, "Unable to release program object");
2899
2900 error = clReleaseProgram(fully_linked_program);
2901 test_error(error, "Unable to release program object");
2902
2903 return 0;
2904 }
2905
test_execute_after_two_file_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2906 int test_execute_after_two_file_link(cl_device_id deviceID, cl_context context,
2907 cl_command_queue queue, int num_elements)
2908 {
2909 int error;
2910 cl_program program, another_program;
2911
2912 log_info("Testing two file compiling and linking and execution of two "
2913 "kernels afterwards ...\n");
2914 error = create_single_kernel_helper_create_program(context, &program, 1,
2915 &simple_kernel);
2916 if (program == NULL || error != CL_SUCCESS)
2917 {
2918 log_error(
2919 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2920 IGetErrorString(error), __FILE__, __LINE__);
2921 return -1;
2922 }
2923
2924 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
2925 NULL);
2926 test_error(error, "Unable to compile a simple program");
2927
2928 error = create_single_kernel_helper_create_program(
2929 context, &another_program, 1, &another_simple_kernel);
2930 if (another_program == NULL || error != CL_SUCCESS)
2931 {
2932 log_error(
2933 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2934 IGetErrorString(error), __FILE__, __LINE__);
2935 return -1;
2936 }
2937
2938 error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL,
2939 NULL, NULL);
2940 test_error(error, "Unable to compile a simple program");
2941
2942 cl_program two_programs[2] = { program, another_program };
2943 cl_program fully_linked_program = clLinkProgram(
2944 context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
2945 test_error(error, "Unable to create an executable from two binaries");
2946
2947 cl_kernel kernel =
2948 clCreateKernel(fully_linked_program, "CopyBuffer", &error);
2949 test_error(error, "Unable to create a simple kernel");
2950
2951 error = verifyCopyBuffer(context, queue, kernel);
2952 if (error != CL_SUCCESS) return error;
2953
2954 cl_kernel another_kernel =
2955 clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
2956 test_error(error, "Unable to create another simple kernel");
2957
2958 error = verifyCopyBuffer(context, queue, another_kernel);
2959 if (error != CL_SUCCESS) return error;
2960
2961 /* All done! */
2962 error = clReleaseKernel(kernel);
2963 test_error(error, "Unable to release kernel object");
2964
2965 error = clReleaseKernel(another_kernel);
2966 test_error(error, "Unable to release another kernel object");
2967
2968 error = clReleaseProgram(program);
2969 test_error(error, "Unable to release program object");
2970
2971 error = clReleaseProgram(another_program);
2972 test_error(error, "Unable to release program object");
2973
2974 error = clReleaseProgram(fully_linked_program);
2975 test_error(error, "Unable to release program object");
2976
2977 return 0;
2978 }
2979
test_execute_after_embedded_header_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)2980 int test_execute_after_embedded_header_link(cl_device_id deviceID,
2981 cl_context context,
2982 cl_command_queue queue,
2983 int num_elements)
2984 {
2985 int error;
2986 cl_program program, header, simple_program;
2987
2988 log_info("Testing execution after embedded header link...\n");
2989 // we will test execution after compiling and linking with embedded headers
2990 program = clCreateProgramWithSource(
2991 context, 1, &another_simple_kernel_with_header, NULL, &error);
2992 if (program == NULL || error != CL_SUCCESS)
2993 {
2994 log_error(
2995 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
2996 IGetErrorString(error), __FILE__, __LINE__);
2997 return -1;
2998 }
2999
3000 header =
3001 clCreateProgramWithSource(context, 1, &simple_header, NULL, &error);
3002 if (header == NULL || error != CL_SUCCESS)
3003 {
3004 log_error(
3005 "ERROR: Unable to create a simple header program! (%s in %s:%d)\n",
3006 IGetErrorString(error), __FILE__, __LINE__);
3007 return -1;
3008 }
3009
3010 error = clCompileProgram(program, 1, &deviceID, NULL, 1, &header,
3011 &simple_header_name, NULL, NULL);
3012 test_error(error,
3013 "Unable to compile a simple program with embedded header");
3014
3015 simple_program =
3016 clCreateProgramWithSource(context, 1, &simple_kernel, NULL, &error);
3017 if (simple_program == NULL || error != CL_SUCCESS)
3018 {
3019 log_error(
3020 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
3021 IGetErrorString(error), __FILE__, __LINE__);
3022 return -1;
3023 }
3024
3025 error = clCompileProgram(simple_program, 1, &deviceID, NULL, 0, NULL, NULL,
3026 NULL, NULL);
3027 test_error(error, "Unable to compile a simple program");
3028
3029 cl_program two_programs[2] = { program, simple_program };
3030 cl_program fully_linked_program = clLinkProgram(
3031 context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
3032 test_error(error,
3033 "Unable to create an executable from two binaries, one compiled "
3034 "with embedded header");
3035
3036 cl_kernel kernel =
3037 clCreateKernel(fully_linked_program, "CopyBuffer", &error);
3038 test_error(error, "Unable to create a simple kernel");
3039
3040 error = verifyCopyBuffer(context, queue, kernel);
3041 if (error != CL_SUCCESS) return error;
3042
3043 cl_kernel another_kernel =
3044 clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
3045 test_error(error, "Unable to create another simple kernel");
3046
3047 error = verifyCopyBuffer(context, queue, another_kernel);
3048 if (error != CL_SUCCESS) return error;
3049
3050 /* All done! */
3051 error = clReleaseKernel(kernel);
3052 test_error(error, "Unable to release kernel object");
3053
3054 error = clReleaseKernel(another_kernel);
3055 test_error(error, "Unable to release another kernel object");
3056
3057 error = clReleaseProgram(program);
3058 test_error(error, "Unable to release program object");
3059
3060 error = clReleaseProgram(header);
3061 test_error(error, "Unable to release program object");
3062
3063 error = clReleaseProgram(simple_program);
3064 test_error(error, "Unable to release program object");
3065
3066 error = clReleaseProgram(fully_linked_program);
3067 test_error(error, "Unable to release program object");
3068
3069 return 0;
3070 }
3071
3072 #if defined(__APPLE__) || defined(__linux)
3073 #define _mkdir(x) mkdir(x, S_IRWXU)
3074 #define _chdir chdir
3075 #define _rmdir rmdir
3076 #define _unlink unlink
3077 #else
3078 #include <direct.h>
3079 #endif
3080
test_execute_after_included_header_link(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)3081 int test_execute_after_included_header_link(cl_device_id deviceID,
3082 cl_context context,
3083 cl_command_queue queue,
3084 int num_elements)
3085 {
3086 int error;
3087 cl_program program, simple_program;
3088
3089 log_info("Testing execution after included header link...\n");
3090 // we will test execution after compiling and linking with included headers
3091 program = clCreateProgramWithSource(
3092 context, 1, &another_simple_kernel_with_header, NULL, &error);
3093 if (program == NULL || error != CL_SUCCESS)
3094 {
3095 log_error(
3096 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
3097 IGetErrorString(error), __FILE__, __LINE__);
3098 return -1;
3099 }
3100
3101 /* setup */
3102 #if (defined(__linux__) || defined(__APPLE__)) && (!defined(__ANDROID__))
3103 /* Some tests systems doesn't allow one to write in the test directory */
3104 if (_chdir("/tmp") != 0)
3105 {
3106 log_error("ERROR: Unable to remove directory foo/bar! (in %s:%d)\n",
3107 __FILE__, __LINE__);
3108 return -1;
3109 }
3110 #endif
3111 if (_mkdir("foo") != 0)
3112 {
3113 log_error("ERROR: Unable to create directory foo! (in %s:%d)\n",
3114 __FILE__, __LINE__);
3115 return -1;
3116 }
3117 if (_mkdir("foo/bar") != 0)
3118 {
3119 log_error("ERROR: Unable to create directory foo/bar! (in %s:%d)\n",
3120 __FILE__, __LINE__);
3121 return -1;
3122 }
3123 if (_chdir("foo/bar") != 0)
3124 {
3125 log_error("ERROR: Unable to change to directory foo/bar! (in %s:%d)\n",
3126 __FILE__, __LINE__);
3127 return -1;
3128 }
3129 FILE *simple_header_file = fopen(simple_header_name, "w");
3130 if (simple_header_file == NULL)
3131 {
3132 log_error("ERROR: Unable to create simple header file %s! (in %s:%d)\n",
3133 simple_header_name, __FILE__, __LINE__);
3134 return -1;
3135 }
3136 if (fprintf(simple_header_file, "%s", simple_header) < 0)
3137 {
3138 log_error(
3139 "ERROR: Unable to write to simple header file %s! (in %s:%d)\n",
3140 simple_header_name, __FILE__, __LINE__);
3141 return -1;
3142 }
3143 if (fclose(simple_header_file) != 0)
3144 {
3145 log_error("ERROR: Unable to close simple header file %s! (in %s:%d)\n",
3146 simple_header_name, __FILE__, __LINE__);
3147 return -1;
3148 }
3149 if (_chdir("../..") != 0)
3150 {
3151 log_error("ERROR: Unable to change to original working directory! (in "
3152 "%s:%d)\n",
3153 __FILE__, __LINE__);
3154 return -1;
3155 }
3156 #if (defined(__linux__) || defined(__APPLE__)) && (!defined(__ANDROID__))
3157 error = clCompileProgram(program, 1, &deviceID, "-I/tmp/foo/bar", 0, NULL,
3158 NULL, NULL, NULL);
3159 #else
3160 error = clCompileProgram(program, 1, &deviceID, "-Ifoo/bar", 0, NULL, NULL,
3161 NULL, NULL);
3162 #endif
3163 test_error(error,
3164 "Unable to compile a simple program with included header");
3165
3166 /* cleanup */
3167 if (_chdir("foo/bar") != 0)
3168 {
3169 log_error("ERROR: Unable to change to directory foo/bar! (in %s:%d)\n",
3170 __FILE__, __LINE__);
3171 return -1;
3172 }
3173 if (_unlink(simple_header_name) != 0)
3174 {
3175 log_error("ERROR: Unable to remove simple header file %s! (in %s:%d)\n",
3176 simple_header_name, __FILE__, __LINE__);
3177 return -1;
3178 }
3179 if (_chdir("../..") != 0)
3180 {
3181 log_error("ERROR: Unable to change to original working directory! (in "
3182 "%s:%d)\n",
3183 __FILE__, __LINE__);
3184 return -1;
3185 }
3186 if (_rmdir("foo/bar") != 0)
3187 {
3188 log_error("ERROR: Unable to remove directory foo/bar! (in %s:%d)\n",
3189 __FILE__, __LINE__);
3190 return -1;
3191 }
3192 if (_rmdir("foo") != 0)
3193 {
3194 log_error("ERROR: Unable to remove directory foo! (in %s:%d)\n",
3195 __FILE__, __LINE__);
3196 return -1;
3197 }
3198
3199 simple_program =
3200 clCreateProgramWithSource(context, 1, &simple_kernel, NULL, &error);
3201 if (simple_program == NULL || error != CL_SUCCESS)
3202 {
3203 log_error(
3204 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
3205 IGetErrorString(error), __FILE__, __LINE__);
3206 return -1;
3207 }
3208
3209 error = clCompileProgram(simple_program, 1, &deviceID, NULL, 0, NULL, NULL,
3210 NULL, NULL);
3211 test_error(error, "Unable to compile a simple program");
3212
3213 cl_program two_programs[2] = { program, simple_program };
3214 cl_program fully_linked_program = clLinkProgram(
3215 context, 1, &deviceID, "", 2, two_programs, NULL, NULL, &error);
3216 test_error(error,
3217 "Unable to create an executable from two binaries, one compiled "
3218 "with embedded header");
3219
3220 cl_kernel kernel =
3221 clCreateKernel(fully_linked_program, "CopyBuffer", &error);
3222 test_error(error, "Unable to create a simple kernel");
3223
3224 error = verifyCopyBuffer(context, queue, kernel);
3225 if (error != CL_SUCCESS) return error;
3226
3227 cl_kernel another_kernel =
3228 clCreateKernel(fully_linked_program, "AnotherCopyBuffer", &error);
3229 test_error(error, "Unable to create another simple kernel");
3230
3231 error = verifyCopyBuffer(context, queue, another_kernel);
3232 if (error != CL_SUCCESS) return error;
3233
3234 /* All done! */
3235 error = clReleaseKernel(kernel);
3236 test_error(error, "Unable to release kernel object");
3237
3238 error = clReleaseKernel(another_kernel);
3239 test_error(error, "Unable to release another kernel object");
3240
3241 error = clReleaseProgram(program);
3242 test_error(error, "Unable to release program object");
3243
3244 error = clReleaseProgram(simple_program);
3245 test_error(error, "Unable to release program object");
3246
3247 error = clReleaseProgram(fully_linked_program);
3248 test_error(error, "Unable to release program object");
3249
3250 return 0;
3251 }
3252
test_program_binary_type(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)3253 int test_program_binary_type(cl_device_id deviceID, cl_context context,
3254 cl_command_queue queue, int num_elements)
3255 {
3256 int error;
3257 cl_program program, another_program, program_with_binary,
3258 fully_linked_program_with_binary;
3259 cl_program_binary_type program_type = -1;
3260 size_t size;
3261 size_t binarySize;
3262 unsigned char *binary;
3263
3264 log_info("Testing querying of program binary type...\n");
3265 error = create_single_kernel_helper_create_program(context, &program, 1,
3266 &simple_kernel);
3267 if (program == NULL || error != CL_SUCCESS)
3268 {
3269 log_error(
3270 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
3271 IGetErrorString(error), __FILE__, __LINE__);
3272 return -1;
3273 }
3274
3275 error = clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL, NULL,
3276 NULL);
3277 test_error(error, "Unable to compile a simple program");
3278
3279 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BINARY_TYPE,
3280 sizeof(cl_program_binary_type), &program_type,
3281 NULL);
3282 test_error(error, "Unable to get program binary type");
3283 if (program_type != CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT)
3284 {
3285 log_error("ERROR: Expected program type of a just compiled program to "
3286 "be CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT (in %s:%d)\n",
3287 __FILE__, __LINE__);
3288 return -1;
3289 }
3290 program_type = -1;
3291
3292 // Get the size of the resulting binary (only one device)
3293 error = clGetProgramInfo(program, CL_PROGRAM_BINARY_SIZES,
3294 sizeof(binarySize), &binarySize, NULL);
3295 test_error(error, "Unable to get binary size");
3296
3297 // Sanity check
3298 if (binarySize == 0)
3299 {
3300 log_error("ERROR: Binary size of program is zero (in %s:%d)\n",
3301 __FILE__, __LINE__);
3302 return -1;
3303 }
3304
3305 // Create a buffer and get the actual binary
3306 {
3307 binary = (unsigned char *)malloc(sizeof(unsigned char) * binarySize);
3308 if (binary == NULL)
3309 {
3310 log_error("ERROR: Unable to allocate binary character array with "
3311 "%lu characters! (in %s:%d)\n",
3312 binarySize, __FILE__, __LINE__);
3313 return -1;
3314 }
3315 unsigned char *buffers[1] = { binary };
3316 cl_int loadErrors[1];
3317
3318 // Do another sanity check here first
3319 size_t size;
3320 error = clGetProgramInfo(program, CL_PROGRAM_BINARIES, 0, NULL, &size);
3321 test_error(error, "Unable to get expected size of binaries array");
3322 if (size != sizeof(buffers))
3323 {
3324 log_error(
3325 "ERROR: Expected size of binaries array in clGetProgramInfo is "
3326 "incorrect (should be %d, got %d) (in %s:%d)\n",
3327 (int)sizeof(buffers), (int)size, __FILE__, __LINE__);
3328 free(binary);
3329 return -1;
3330 }
3331
3332 error = clGetProgramInfo(program, CL_PROGRAM_BINARIES, sizeof(buffers),
3333 &buffers, NULL);
3334 test_error(error, "Unable to get program binary");
3335
3336 // use clCreateProgramWithBinary
3337 program_with_binary = clCreateProgramWithBinary(
3338 context, 1, &deviceID, &binarySize, (const unsigned char **)buffers,
3339 loadErrors, &error);
3340 test_error(error, "Unable to create program with binary");
3341
3342 error = clGetProgramBuildInfo(
3343 program_with_binary, deviceID, CL_PROGRAM_BINARY_TYPE,
3344 sizeof(cl_program_binary_type), &program_type, NULL);
3345 test_error(error, "Unable to get program binary type");
3346 if (program_type != CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT)
3347 {
3348 log_error("ERROR: Expected program type of a program created from "
3349 "compiled object to be "
3350 "CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT (in %s:%d)\n",
3351 __FILE__, __LINE__);
3352 return -1;
3353 }
3354 program_type = -1;
3355 free(binary);
3356 }
3357
3358 cl_program my_newly_minted_library =
3359 clLinkProgram(context, 1, &deviceID, "-create-library", 1,
3360 &program_with_binary, NULL, NULL, &error);
3361 test_error(error, "Unable to create a simple library");
3362 error = clGetProgramBuildInfo(
3363 my_newly_minted_library, deviceID, CL_PROGRAM_BINARY_TYPE,
3364 sizeof(cl_program_binary_type), &program_type, NULL);
3365 test_error(error, "Unable to get program binary type");
3366 if (program_type != CL_PROGRAM_BINARY_TYPE_LIBRARY)
3367 {
3368 log_error("ERROR: Expected program type of a just linked library to be "
3369 "CL_PROGRAM_BINARY_TYPE_LIBRARY (in %s:%d)\n",
3370 __FILE__, __LINE__);
3371 return -1;
3372 }
3373 program_type = -1;
3374
3375 // Get the size of the resulting library (only one device)
3376 error = clGetProgramInfo(my_newly_minted_library, CL_PROGRAM_BINARY_SIZES,
3377 sizeof(binarySize), &binarySize, NULL);
3378 test_error(error, "Unable to get binary size");
3379
3380 // Sanity check
3381 if (binarySize == 0)
3382 {
3383 log_error("ERROR: Binary size of program is zero (in %s:%d)\n",
3384 __FILE__, __LINE__);
3385 return -1;
3386 }
3387
3388 // Create a buffer and get the actual binary
3389 binary = (unsigned char *)malloc(sizeof(unsigned char) * binarySize);
3390 if (binary == NULL)
3391 {
3392 log_error("ERROR: Unable to allocate binary character array with %lu "
3393 "characters! (in %s:%d)\n",
3394 binarySize, __FILE__, __LINE__);
3395 return -1;
3396 }
3397
3398 unsigned char *buffers[1] = { binary };
3399 cl_int loadErrors[1];
3400
3401 // Do another sanity check here first
3402 error = clGetProgramInfo(my_newly_minted_library, CL_PROGRAM_BINARIES, 0,
3403 NULL, &size);
3404 test_error(error, "Unable to get expected size of binaries array");
3405 if (size != sizeof(buffers))
3406 {
3407 log_error("ERROR: Expected size of binaries array in clGetProgramInfo "
3408 "is incorrect (should be %d, got %d) (in %s:%d)\n",
3409 (int)sizeof(buffers), (int)size, __FILE__, __LINE__);
3410 free(binary);
3411 return -1;
3412 }
3413
3414 error = clGetProgramInfo(my_newly_minted_library, CL_PROGRAM_BINARIES,
3415 sizeof(buffers), &buffers, NULL);
3416 test_error(error, "Unable to get program binary");
3417
3418 // use clCreateProgramWithBinary
3419 cl_program library_with_binary = clCreateProgramWithBinary(
3420 context, 1, &deviceID, &binarySize, (const unsigned char **)buffers,
3421 loadErrors, &error);
3422 test_error(error, "Unable to create program with binary");
3423 error = clGetProgramBuildInfo(
3424 library_with_binary, deviceID, CL_PROGRAM_BINARY_TYPE,
3425 sizeof(cl_program_binary_type), &program_type, NULL);
3426 test_error(error, "Unable to get program binary type");
3427 if (program_type != CL_PROGRAM_BINARY_TYPE_LIBRARY)
3428 {
3429 log_error("ERROR: Expected program type of a library loaded with "
3430 "binary to be CL_PROGRAM_BINARY_TYPE_LIBRARY (in %s:%d)\n",
3431 __FILE__, __LINE__);
3432 return -1;
3433 }
3434 program_type = -1;
3435 free(binary);
3436
3437 error = create_single_kernel_helper_create_program(
3438 context, &another_program, 1, &another_simple_kernel);
3439 if (another_program == NULL || error != CL_SUCCESS)
3440 {
3441 log_error(
3442 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
3443 IGetErrorString(error), __FILE__, __LINE__);
3444 return -1;
3445 }
3446
3447 error = clCompileProgram(another_program, 1, &deviceID, NULL, 0, NULL, NULL,
3448 NULL, NULL);
3449 test_error(error, "Unable to compile a simple program");
3450
3451 cl_program program_and_archive[2] = { another_program,
3452 library_with_binary };
3453 cl_program fully_linked_program = clLinkProgram(
3454 context, 1, &deviceID, "", 2, program_and_archive, NULL, NULL, &error);
3455 test_error(error,
3456 "Unable to create an executable from a binary and a library");
3457
3458 error = clGetProgramBuildInfo(
3459 fully_linked_program, deviceID, CL_PROGRAM_BINARY_TYPE,
3460 sizeof(cl_program_binary_type), &program_type, NULL);
3461 test_error(error, "Unable to get program binary type");
3462 if (program_type != CL_PROGRAM_BINARY_TYPE_EXECUTABLE)
3463 {
3464 log_error("ERROR: Expected program type of a newly build executable to "
3465 "be CL_PROGRAM_BINARY_TYPE_EXECUTABLE (in %s:%d)\n",
3466 __FILE__, __LINE__);
3467 return -1;
3468 }
3469 program_type = -1;
3470
3471 // Get the size of the resulting binary (only one device)
3472 error = clGetProgramInfo(fully_linked_program, CL_PROGRAM_BINARY_SIZES,
3473 sizeof(binarySize), &binarySize, NULL);
3474 test_error(error, "Unable to get binary size");
3475
3476 // Sanity check
3477 if (binarySize == 0)
3478 {
3479 log_error("ERROR: Binary size of program is zero (in %s:%d)\n",
3480 __FILE__, __LINE__);
3481 return -1;
3482 }
3483
3484 // Create a buffer and get the actual binary
3485 {
3486 binary = (unsigned char *)malloc(sizeof(unsigned char) * binarySize);
3487 if (binary == NULL)
3488 {
3489 log_error("ERROR: Unable to allocate binary character array with "
3490 "%lu characters! (in %s:%d)\n",
3491 binarySize, __FILE__, __LINE__);
3492 return -1;
3493 }
3494 unsigned char *buffers[1] = { binary };
3495 cl_int loadErrors[1];
3496
3497 // Do another sanity check here first
3498 size_t size;
3499 error = clGetProgramInfo(fully_linked_program, CL_PROGRAM_BINARIES, 0,
3500 NULL, &size);
3501 test_error(error, "Unable to get expected size of binaries array");
3502 if (size != sizeof(buffers))
3503 {
3504 log_error(
3505 "ERROR: Expected size of binaries array in clGetProgramInfo is "
3506 "incorrect (should be %d, got %d) (in %s:%d)\n",
3507 (int)sizeof(buffers), (int)size, __FILE__, __LINE__);
3508 free(binary);
3509 return -1;
3510 }
3511
3512 error = clGetProgramInfo(fully_linked_program, CL_PROGRAM_BINARIES,
3513 sizeof(buffers), &buffers, NULL);
3514 test_error(error, "Unable to get program binary");
3515
3516 // use clCreateProgramWithBinary
3517 fully_linked_program_with_binary = clCreateProgramWithBinary(
3518 context, 1, &deviceID, &binarySize, (const unsigned char **)buffers,
3519 loadErrors, &error);
3520 test_error(error, "Unable to create program with binary");
3521
3522 error = clGetProgramBuildInfo(
3523 fully_linked_program_with_binary, deviceID, CL_PROGRAM_BINARY_TYPE,
3524 sizeof(cl_program_binary_type), &program_type, NULL);
3525 test_error(error, "Unable to get program binary type");
3526 if (program_type != CL_PROGRAM_BINARY_TYPE_EXECUTABLE)
3527 {
3528 log_error("ERROR: Expected program type of a program created from "
3529 "a fully linked executable binary to be "
3530 "CL_PROGRAM_BINARY_TYPE_EXECUTABLE (in %s:%d)\n",
3531 __FILE__, __LINE__);
3532 return -1;
3533 }
3534 program_type = -1;
3535 free(binary);
3536 }
3537
3538 error = clBuildProgram(fully_linked_program_with_binary, 1, &deviceID, NULL,
3539 NULL, NULL);
3540 test_error(error, "Unable to build a simple program");
3541
3542 cl_kernel kernel =
3543 clCreateKernel(fully_linked_program_with_binary, "CopyBuffer", &error);
3544 test_error(error, "Unable to create a simple kernel");
3545
3546 error = verifyCopyBuffer(context, queue, kernel);
3547 if (error != CL_SUCCESS) return error;
3548
3549 cl_kernel another_kernel = clCreateKernel(fully_linked_program_with_binary,
3550 "AnotherCopyBuffer", &error);
3551 test_error(error, "Unable to create another simple kernel");
3552
3553 error = verifyCopyBuffer(context, queue, another_kernel);
3554 if (error != CL_SUCCESS) return error;
3555
3556 /* All done! */
3557 error = clReleaseKernel(kernel);
3558 test_error(error, "Unable to release kernel object");
3559
3560 error = clReleaseKernel(another_kernel);
3561 test_error(error, "Unable to release another kernel object");
3562
3563 error = clReleaseProgram(program);
3564 test_error(error, "Unable to release program object");
3565
3566 /* Oh, one more thing. Steve Jobs and apparently Herb Sutter. The question
3567 * is "Who is copying whom?" */
3568 error = create_single_kernel_helper_create_program(context, &program, 1,
3569 &simple_kernel);
3570 if (program == NULL || error != CL_SUCCESS)
3571 {
3572 log_error(
3573 "ERROR: Unable to create a simple test program! (%s in %s:%d)\n",
3574 IGetErrorString(error), __FILE__, __LINE__);
3575 return -1;
3576 }
3577
3578 error = clBuildProgram(program, 1, &deviceID, NULL, NULL, NULL);
3579 test_error(error, "Unable to build a simple program");
3580 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BINARY_TYPE,
3581 sizeof(cl_program_binary_type), &program_type,
3582 NULL);
3583 test_error(error, "Unable to get program binary type");
3584 if (program_type != CL_PROGRAM_BINARY_TYPE_EXECUTABLE)
3585 {
3586 log_error(
3587 "ERROR: Expected program type of a program created from compiled "
3588 "object to be CL_PROGRAM_BINARY_TYPE_EXECUTABLE (in %s:%d)\n",
3589 __FILE__, __LINE__);
3590 return -1;
3591 }
3592 program_type = -1;
3593
3594 /* All's well that ends well. William Shakespeare */
3595 error = clReleaseProgram(program);
3596 test_error(error, "Unable to release program object");
3597
3598 error = clReleaseProgram(another_program);
3599 test_error(error, "Unable to release program object");
3600
3601 error = clReleaseProgram(my_newly_minted_library);
3602 test_error(error, "Unable to release program object");
3603
3604 error = clReleaseProgram(library_with_binary);
3605 test_error(error, "Unable to release program object");
3606
3607 error = clReleaseProgram(fully_linked_program);
3608 test_error(error, "Unable to release program object");
3609
3610 error = clReleaseProgram(fully_linked_program_with_binary);
3611 test_error(error, "Unable to release program object");
3612
3613 error = clReleaseProgram(program_with_binary);
3614 test_error(error, "Unable to release program object");
3615
3616 return 0;
3617 }
3618
3619 volatile int compileNotificationSent;
3620
test_notify_compile_complete(cl_program program,void * userData)3621 void CL_CALLBACK test_notify_compile_complete(cl_program program,
3622 void *userData)
3623 {
3624 if (userData == NULL || strcmp((char *)userData, "compilation") != 0)
3625 {
3626 log_error("ERROR: User data passed in to compile notify function was "
3627 "not correct! (in %s:%d)\n",
3628 __FILE__, __LINE__);
3629 compileNotificationSent = -1;
3630 }
3631 else
3632 compileNotificationSent = 1;
3633 log_info("\n <-- program successfully compiled\n");
3634 }
3635
3636 volatile int libraryCreationNotificationSent;
3637
test_notify_create_library_complete(cl_program program,void * userData)3638 void CL_CALLBACK test_notify_create_library_complete(cl_program program,
3639 void *userData)
3640 {
3641 if (userData == NULL || strcmp((char *)userData, "create library") != 0)
3642 {
3643 log_error("ERROR: User data passed in to library creation notify "
3644 "function was not correct! (in %s:%d)\n",
3645 __FILE__, __LINE__);
3646 libraryCreationNotificationSent = -1;
3647 }
3648 else
3649 libraryCreationNotificationSent = 1;
3650 log_info("\n <-- library successfully created\n");
3651 }
3652
3653 volatile int linkNotificationSent;
3654
test_notify_link_complete(cl_program program,void * userData)3655 void CL_CALLBACK test_notify_link_complete(cl_program program, void *userData)
3656 {
3657 if (userData == NULL || strcmp((char *)userData, "linking") != 0)
3658 {
3659 log_error("ERROR: User data passed in to link notify function was not "
3660 "correct! (in %s:%d)\n",
3661 __FILE__, __LINE__);
3662 linkNotificationSent = -1;
3663 }
3664 else
3665 linkNotificationSent = 1;
3666 log_info("\n <-- program successfully linked\n");
3667 }
3668
test_large_compile_and_link_status_options_log(cl_context context,cl_device_id deviceID,cl_command_queue queue,unsigned int numLines)3669 int test_large_compile_and_link_status_options_log(cl_context context,
3670 cl_device_id deviceID,
3671 cl_command_queue queue,
3672 unsigned int numLines)
3673 {
3674 int error;
3675 cl_program program;
3676 cl_program *simple_kernels;
3677 const char **lines;
3678 unsigned int i;
3679 char buffer[MAX_LINE_SIZE_IN_PROGRAM];
3680 char *compile_log;
3681 char *compile_options;
3682 char *library_log;
3683 char *library_options;
3684 char *linking_log;
3685 char *linking_options;
3686 cl_build_status status;
3687 size_t size_ret;
3688
3689 compileNotificationSent = libraryCreationNotificationSent =
3690 linkNotificationSent = 0;
3691
3692 simple_kernels = (cl_program *)malloc(numLines * sizeof(cl_program));
3693 if (simple_kernels == NULL)
3694 {
3695 log_error("ERROR: Unable to allocate kernels array with %d kernels! "
3696 "(in %s:%d)\n",
3697 numLines, __FILE__, __LINE__);
3698 return -1;
3699 }
3700 /* First, allocate the array for our line pointers */
3701 lines = (const char **)malloc((2 * numLines + 2) * sizeof(const char *));
3702 if (lines == NULL)
3703 {
3704 log_error(
3705 "ERROR: Unable to allocate lines array with %d lines! (in %s:%d)\n",
3706 (2 * numLines + 2), __FILE__, __LINE__);
3707 return -1;
3708 }
3709
3710 for (i = 0; i < numLines; i++)
3711 {
3712 sprintf(buffer, composite_kernel_extern_template, i);
3713 lines[i] = _strdup(buffer);
3714 }
3715 /* First and last lines are easy */
3716 lines[numLines] = composite_kernel_start;
3717 lines[2 * numLines + 1] = composite_kernel_end;
3718
3719 /* Fill the rest with templated kernels */
3720 for (i = numLines + 1; i < 2 * numLines + 1; i++)
3721 {
3722 sprintf(buffer, composite_kernel_template, i - numLines - 1);
3723 lines[i] = _strdup(buffer);
3724 }
3725
3726 /* Try to create a program with these lines */
3727 error = create_single_kernel_helper_create_program(context, &program,
3728 2 * numLines + 2, lines);
3729 if (program == NULL || error != CL_SUCCESS)
3730 {
3731 log_error("ERROR: Unable to create long test program with %d lines! "
3732 "(%s) (in %s:%d)\n",
3733 numLines, IGetErrorString(error), __FILE__, __LINE__);
3734 return -1;
3735 }
3736
3737 /* Lets check that the compilation status is CL_BUILD_NONE */
3738 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS,
3739 sizeof(status), &status, NULL);
3740 test_error(error, "Unable to get program compile status");
3741 if (status != CL_BUILD_NONE)
3742 {
3743 log_error("ERROR: Expected compile status to be CL_BUILD_NONE prior to "
3744 "the beginning of the compilation! (status: %d in %s:%d)\n",
3745 (int)status, __FILE__, __LINE__);
3746 return -1;
3747 }
3748
3749 /* Compile it */
3750 error =
3751 clCompileProgram(program, 1, &deviceID, NULL, 0, NULL, NULL,
3752 test_notify_compile_complete, (void *)"compilation");
3753 test_error(error, "Unable to compile a simple program");
3754
3755 /* Wait for compile to complete (just keep polling, since we're just a test
3756 */
3757 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS,
3758 sizeof(status), &status, NULL);
3759 test_error(error, "Unable to get program compile status");
3760
3761 while ((int)status == CL_BUILD_IN_PROGRESS)
3762 {
3763 log_info("\n -- still waiting for compile... (status is %d)", status);
3764 sleep(1);
3765 error =
3766 clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_STATUS,
3767 sizeof(status), &status, NULL);
3768 test_error(error, "Unable to get program compile status");
3769 }
3770 if (status != CL_BUILD_SUCCESS)
3771 {
3772 log_error("ERROR: compile failed! (status: %d in %s:%d)\n", (int)status,
3773 __FILE__, __LINE__);
3774 return -1;
3775 }
3776
3777 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG, 0,
3778 NULL, &size_ret);
3779 test_error(error, "Device failed to return compile log size");
3780 compile_log = (char *)malloc(size_ret);
3781 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_LOG,
3782 size_ret, compile_log, NULL);
3783 if (error != CL_SUCCESS)
3784 {
3785 log_error("Device failed to return a compile log (in %s:%d)\n",
3786 __FILE__, __LINE__);
3787 test_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed");
3788 }
3789 log_info("BUILD LOG: %s\n", compile_log);
3790 free(compile_log);
3791
3792 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_OPTIONS,
3793 0, NULL, &size_ret);
3794 test_error(error, "Device failed to return compile options size");
3795 compile_options = (char *)malloc(size_ret);
3796 error = clGetProgramBuildInfo(program, deviceID, CL_PROGRAM_BUILD_OPTIONS,
3797 size_ret, compile_options, NULL);
3798 test_error(
3799 error,
3800 "Device failed to return compile options.\nclGetProgramBuildInfo "
3801 "CL_PROGRAM_BUILD_OPTIONS failed");
3802
3803 log_info("BUILD OPTIONS: %s\n", compile_options);
3804 free(compile_options);
3805
3806 /* Create and compile templated kernels */
3807 for (i = 0; i < numLines; i++)
3808 {
3809 sprintf(buffer, simple_kernel_template, i);
3810 const char *kernel_source = _strdup(buffer);
3811 error = create_single_kernel_helper_create_program(
3812 context, &simple_kernels[i], 1, &kernel_source);
3813 if (simple_kernels[i] == NULL || error != CL_SUCCESS)
3814 {
3815 log_error("ERROR: Unable to create long test program with %d "
3816 "lines! (%s in %s:%d)",
3817 numLines, IGetErrorString(error), __FILE__, __LINE__);
3818 return -1;
3819 }
3820
3821 /* Compile it */
3822 error = clCompileProgram(simple_kernels[i], 1, &deviceID, NULL, 0, NULL,
3823 NULL, NULL, NULL);
3824 test_error(error, "Unable to compile a simple program");
3825
3826 free((void *)kernel_source);
3827 }
3828
3829 /* Create library out of compiled templated kernels */
3830 cl_program my_newly_minted_library = clLinkProgram(
3831 context, 1, &deviceID, "-create-library", numLines, simple_kernels,
3832 test_notify_create_library_complete, (void *)"create library", &error);
3833 test_error(error, "Unable to create a multi-line library");
3834
3835 /* Wait for library creation to complete (just keep polling, since we're
3836 * just a test */
3837 error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
3838 CL_PROGRAM_BUILD_STATUS, sizeof(status),
3839 &status, NULL);
3840 test_error(error, "Unable to get library creation link status");
3841
3842 while ((int)status == CL_BUILD_IN_PROGRESS)
3843 {
3844 log_info("\n -- still waiting for library creation... (status is %d)",
3845 status);
3846 sleep(1);
3847 error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
3848 CL_PROGRAM_BUILD_STATUS, sizeof(status),
3849 &status, NULL);
3850 test_error(error, "Unable to get library creation link status");
3851 }
3852 if (status != CL_BUILD_SUCCESS)
3853 {
3854 log_error("ERROR: library creation failed! (status: %d in %s:%d)\n",
3855 (int)status, __FILE__, __LINE__);
3856 return -1;
3857 }
3858 error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
3859 CL_PROGRAM_BUILD_LOG, 0, NULL, &size_ret);
3860 test_error(error, "Device failed to return a library creation log size");
3861 library_log = (char *)malloc(size_ret);
3862 error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
3863 CL_PROGRAM_BUILD_LOG, size_ret, library_log,
3864 NULL);
3865 if (error != CL_SUCCESS)
3866 {
3867 log_error("Device failed to return a library creation log (in %s:%d)\n",
3868 __FILE__, __LINE__);
3869 test_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed");
3870 }
3871 log_info("CREATE LIBRARY LOG: %s\n", library_log);
3872 free(library_log);
3873
3874 error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
3875 CL_PROGRAM_BUILD_OPTIONS, 0, NULL, &size_ret);
3876 test_error(error, "Device failed to return library creation options size");
3877 library_options = (char *)malloc(size_ret);
3878 error = clGetProgramBuildInfo(my_newly_minted_library, deviceID,
3879 CL_PROGRAM_BUILD_OPTIONS, size_ret,
3880 library_options, NULL);
3881 test_error(
3882 error,
3883 "Device failed to return library creation "
3884 "options.\nclGetProgramBuildInfo CL_PROGRAM_BUILD_OPTIONS failed");
3885
3886 log_info("CREATE LIBRARY OPTIONS: %s\n", library_options);
3887 free(library_options);
3888
3889 /* Link the program that calls the kernels and the library that contains
3890 * them */
3891 cl_program programs[2] = { program, my_newly_minted_library };
3892 cl_program my_newly_linked_program =
3893 clLinkProgram(context, 1, &deviceID, NULL, 2, programs,
3894 test_notify_link_complete, (void *)"linking", &error);
3895 test_error(error, "Unable to link a program with a library");
3896
3897 /* Wait for linking to complete (just keep polling, since we're just a test
3898 */
3899 error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
3900 CL_PROGRAM_BUILD_STATUS, sizeof(status),
3901 &status, NULL);
3902 test_error(error, "Unable to get program link status");
3903
3904 while ((int)status == CL_BUILD_IN_PROGRESS)
3905 {
3906 log_info("\n -- still waiting for program linking... (status is %d)",
3907 status);
3908 sleep(1);
3909 error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
3910 CL_PROGRAM_BUILD_STATUS, sizeof(status),
3911 &status, NULL);
3912 test_error(error, "Unable to get program link status");
3913 }
3914 if (status != CL_BUILD_SUCCESS)
3915 {
3916 log_error("ERROR: program linking failed! (status: %d in %s:%d)\n",
3917 (int)status, __FILE__, __LINE__);
3918 return -1;
3919 }
3920 error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
3921 CL_PROGRAM_BUILD_LOG, 0, NULL, &size_ret);
3922 test_error(error, "Device failed to return a linking log size");
3923 linking_log = (char *)malloc(size_ret);
3924 error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
3925 CL_PROGRAM_BUILD_LOG, size_ret, linking_log,
3926 NULL);
3927 if (error != CL_SUCCESS)
3928 {
3929 log_error("Device failed to return a linking log (in %s:%d).\n",
3930 __FILE__, __LINE__);
3931 test_error(error, "clGetProgramBuildInfo CL_PROGRAM_BUILD_LOG failed");
3932 }
3933 log_info("BUILDING LOG: %s\n", linking_log);
3934 free(linking_log);
3935
3936 error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
3937 CL_PROGRAM_BUILD_OPTIONS, 0, NULL, &size_ret);
3938 test_error(error, "Device failed to return linking options size");
3939 linking_options = (char *)malloc(size_ret);
3940 error = clGetProgramBuildInfo(my_newly_linked_program, deviceID,
3941 CL_PROGRAM_BUILD_OPTIONS, size_ret,
3942 linking_options, NULL);
3943 test_error(
3944 error,
3945 "Device failed to return linking options.\nclGetProgramBuildInfo "
3946 "CL_PROGRAM_BUILD_OPTIONS failed");
3947
3948 log_info("BUILDING OPTIONS: %s\n", linking_options);
3949 free(linking_options);
3950
3951 // Create the composite kernel
3952 cl_kernel kernel =
3953 clCreateKernel(my_newly_linked_program, "CompositeKernel", &error);
3954 test_error(error, "Unable to create a composite kernel");
3955
3956 // Run the composite kernel and verify the results
3957 error = verifyCopyBuffer(context, queue, kernel);
3958 if (error != CL_SUCCESS) return error;
3959
3960 /* All done! */
3961 error = clReleaseKernel(kernel);
3962 test_error(error, "Unable to release kernel object");
3963
3964 error = clReleaseProgram(program);
3965 test_error(error, "Unable to release program object");
3966
3967 for (i = 0; i < numLines; i++)
3968 {
3969 free((void *)lines[i]);
3970 free((void *)lines[i + numLines + 1]);
3971 }
3972 free(lines);
3973
3974 for (i = 0; i < numLines; i++)
3975 {
3976 error = clReleaseProgram(simple_kernels[i]);
3977 test_error(error, "Unable to release program object");
3978 }
3979 free(simple_kernels);
3980
3981 error = clReleaseProgram(my_newly_minted_library);
3982 test_error(error, "Unable to release program object");
3983
3984 error = clReleaseProgram(my_newly_linked_program);
3985 test_error(error, "Unable to release program object");
3986
3987 return 0;
3988 }
3989
test_compile_and_link_status_options_log(cl_device_id deviceID,cl_context context,cl_command_queue queue,int num_elements)3990 int test_compile_and_link_status_options_log(cl_device_id deviceID,
3991 cl_context context,
3992 cl_command_queue queue,
3993 int num_elements)
3994 {
3995 unsigned int toTest[] = { 256, 0 }; // 512, 1024, 8192, 16384, 32768, 0 };
3996 unsigned int i;
3997
3998 log_info("Testing Compile and Link Status, Options and Logging ...this "
3999 "might take awhile...\n");
4000
4001 for (i = 0; toTest[i] != 0; i++)
4002 {
4003 log_info(" %d...\n", toTest[i]);
4004
4005 #if defined(_WIN32)
4006 clock_t start = clock();
4007 #elif defined(__linux__) || defined(__APPLE__)
4008 timeval time1, time2;
4009 gettimeofday(&time1, NULL);
4010 #endif
4011
4012 if (test_large_compile_and_link_status_options_log(context, deviceID,
4013 queue, toTest[i])
4014 != 0)
4015 {
4016 log_error(
4017 "ERROR: large program compilation, linking, status, options "
4018 "and logging test failed for %d lines! (in %s:%d)\n",
4019 toTest[i], __FILE__, __LINE__);
4020 return -1;
4021 }
4022
4023 #if defined(_WIN32)
4024 clock_t end = clock();
4025 log_perf((float)(end - start) / (float)CLOCKS_PER_SEC, false,
4026 "clock() time in secs", "%d lines", toTest[i]);
4027 #elif defined(__linux__) || defined(__APPLE__)
4028 gettimeofday(&time2, NULL);
4029 log_perf((float)(float)(time2.tv_sec - time1.tv_sec)
4030 + 1.0e-6 * (time2.tv_usec - time1.tv_usec),
4031 false, "wall time in secs", "%d lines", toTest[i]);
4032 #endif
4033 }
4034
4035 return 0;
4036 }
4037